GraphicsMagick 的 OpenCL 开发记录（三十三）

2024-02-08 10:20
文章标签 开发 记录 三十三 opencl graphicsmagick

本文主要是介绍GraphicsMagick 的 OpenCL 开发记录（三十三），希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

文章目录

  • 如何写`ScaleImage()`的硬件加速函数（七）

<2022-04-28 周四>

如何写ScaleImage()的硬件加速函数（七）

其实“如何写ScaleImage()的硬件加速函数（六）”的实现就是一个ResizeHorizontalFilter()将y改成y / xFactor的精简版，并不是ScaleImage()的硬件加速函数。虽然它不是，但至少省掉了ResizeVerticalFilter()的调用，速度上更快了。

但是目前发现的问题还是竖条纹，连续多次缩小一倍，最终图片被黑色竖条纹全部覆盖住，不断缩小或者放大，右侧会出现密集竖条纹，等等等的问题啦。

经过分析，黑色竖纹的产生原因是因为kernel函数ScaleFilter()的最内层的循环没有执行，导致将初始值0.0f赋进了目标地址。

for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++)
{float4 cp = (float4)0.0f;__local CLQuantum* p = inputImageCache + (cacheIndex * 4);cp.x = (float)*(p);cp.y = (float)*(p + 1);cp.z = (float)*(p + 2);cp.w = (float)*(p + 3);filteredPixel += cp;
}

可以这样解决：

STRINGIFY(__kernel __attribute__((reqd_work_group_size(256, 1, 1)))void ScaleFilter(const __global CLQuantum* inputImage, const unsigned int matte_or_cmyk,const unsigned int inputColumns, const unsigned int inputRows, __global CLQuantum* filteredImage,const unsigned int filteredColumns, const unsigned int filteredRows,const float resizeFilterScale,__local CLQuantum* inputImageCache, const int numCachedPixels,const unsigned int pixelPerWorkgroup, const unsigned int pixelChunkSize,__local float4* outputPixelCache, __local float* densityCache, __local float* gammaCache)
{// calculate the range of resized image pixels computed by this workgroupconst unsigned int startX = get_group_id(0) * pixelPerWorkgroup;const unsigned int stopX = MagickMin(startX + pixelPerWorkgroup, filteredColumns);const unsigned int actualNumPixelToCompute = stopX - startX;float xFactor = (float)filteredColumns / inputColumns;// calculate the range of input image pixels to cacheconst int cacheRangeStartX = MagickMax((int)((startX + 0.5f) / xFactor), (int)(0));const int cacheRangeEndX = MagickMin((int)(cacheRangeStartX + numCachedPixels), (int)inputColumns);// cache the input pixels into local memoryconst unsigned int y = get_global_id(1);const unsigned int pos = getPixelIndex(4, inputColumns, cacheRangeStartX, y / xFactor);const unsigned int num_elements = (cacheRangeEndX - cacheRangeStartX) * 4;event_t e = async_work_group_copy(inputImageCache, inputImage + pos, num_elements, 0);wait_group_events(1, &e);unsigned int totalNumChunks = (actualNumPixelToCompute + pixelChunkSize - 1) / pixelChunkSize;for (unsigned int chunk = 0; chunk < totalNumChunks; chunk++){const unsigned int chunkStartX = startX + chunk * pixelChunkSize;const unsigned int chunkStopX = MagickMin(chunkStartX + pixelChunkSize, stopX);const unsigned int actualNumPixelInThisChunk = chunkStopX - chunkStartX;// determine which resized pixel computed by this workitemconst unsigned int itemID = get_local_id(0);const unsigned int numItems = getNumWorkItemsPerPixel(actualNumPixelInThisChunk, get_local_size(0));const int pixelIndex = pixelToCompute(itemID, actualNumPixelInThisChunk, get_local_size(0));float4 filteredPixel = (float4)0.0f;// -1 means this workitem doesn't participate in the computationif (pixelIndex != -1){// x coordinated of the resized pixel computed by this workitemconst int x = chunkStartX + pixelIndex;// calculate how many steps required for this pixelconst float bisect = (x + 0.5) / xFactor + MagickEpsilon;const unsigned int start = (unsigned int)MagickMax(bisect, 0.0f);const unsigned int stop = (unsigned int)MagickMin(bisect + 1, (float)inputColumns);const unsigned int n = stop - start;// calculate how many steps this workitem will contributeunsigned int numStepsPerWorkItem = n / numItems;numStepsPerWorkItem += ((numItems * numStepsPerWorkItem) == n ? 0 : 1);const unsigned int startStep = (itemID % numItems) * numStepsPerWorkItem;if (startStep < n){const unsigned int stopStep = MagickMin(startStep + numStepsPerWorkItem, n);unsigned int cacheIndex = start + startStep - cacheRangeStartX;for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++){float4 cp = (float4)0.0f;__local CLQuantum* p = inputImageCache + (cacheIndex * 4);cp.x = (float)*(p);cp.y = (float)*(p + 1);cp.z = (float)*(p + 2);cp.w = (float)*(p + 3);filteredPixel += cp;}}}if (itemID < actualNumPixelInThisChunk) {outputPixelCache[itemID] = (float4)0.0f;}barrier(CLK_LOCAL_MEM_FENCE);for (unsigned int i = 0; i < numItems; i++) {if (pixelIndex != -1) {if (itemID % numItems == i) {outputPixelCache[pixelIndex] += filteredPixel;}}barrier(CLK_LOCAL_MEM_FENCE);}if (itemID < actualNumPixelInThisChunk){float4 filteredPixel = outputPixelCache[itemID];WriteAllChannels(filteredImage, 4, filteredColumns, chunkStartX + itemID, y, filteredPixel);}}
}
)

测试了一下性能，感觉提升不少（原图缩小一半，共三次操作，原图连续放大一倍两次，共三次操作）：

ScaleImage()加速版本：

20220428104719 0:3.229821  1.672 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104719 0:3.230185  1.672 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 1360
20220428104725 0:9.628057  1.875 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104725 0:9.628288  1.875 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 0
20220428104732 0:16.078872 2.234 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104732 0:16.079057 2.234 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 0
20220428104740 0:24.253815 2.484 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104740 0:24.254118 2.484 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 0
20220428104749 0:33.888819 2.875 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104749 0:33.889007 2.875 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 31
20220428104752 0:36.173104 3.047 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104752 0:36.173301 3.047 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 156
20220428104800 0:44.287153 3.469 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104800 0:44.287372 3.469 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 47
20220428104801 0:45.546271 3.656 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104801 0:45.546588 3.656 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 140
20220428104806 0:49.973027 4.047 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104806 0:49.973217 4.047 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 31
20220428104806 0:50.640522 4.250 11552 opencl.c AcquireOpenCLKernel 744 Accelerate Event Using kernel: ScaleFilter
20220428104806 0:50.640730 4.250 11552 resize.c ScaleImage 1764 Accelerate Event accelerate scale: 141

ScaleImage()原先版本：

20220428104934 0:1.982873  0.266 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104934 0:2.040677  0.328 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 63
20220428104940 0:7.854823  0.578 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104940 0:7.913365  0.625 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 47
20220428104944 0:11.896725 0.875 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104944 0:11.956722 0.938 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 63
20220428104951 0:18.070817 1.219 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104951 0:18.378405 1.516 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 297
20220428104952 0:19.394056 1.531 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104953 0:20.634341 2.781 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 1250
20220428104958 0:25.534006 3.063 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428104958 0:25.836584 3.375 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 312
20220428104959 0:26.729520 3.406 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428105000 0:27.930533 4.609 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 1203
20220428105011 0:38.879392 5.438 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428105012 0:39.210382 5.766 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 328
20220428105012 0:39.872525 5.797 10052 resize.c ScaleImage 1770 Accelerate Event AccelerateScaleImage null
20220428105014 0:41.176969 7.094 10052 resize.c ScaleImage 2116 Accelerate Event normal scale: 1297

这篇关于GraphicsMagick 的 OpenCL 开发记录（三十三）的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/690726。

相关文章

这15个Vue指令，让你的项目开发爽到爆

这15个Vue指令，让你的项目开发爽到爆

1. V-Hotkey 仓库地址: github.com/Dafrok/v-ho… Demo: 戳这里 https://dafrok.github.io/v-hotkey 安装: npm install --save v-hotkey 这个指令可以给组件绑定一个或多个快捷键。你想要通过按下 Escape 键后隐藏某个组件，按住 Control 和回车键再显示它吗？小菜一碟： <template
阅读更多...
Hadoop企业开发案例调优场景

Hadoop企业开发案例调优场景

需求 （1）需求：从1G数据中，统计每个单词出现次数。服务器3台，每台配置4G内存，4核CPU，4线程。 （2）需求分析： 1G / 128m = 8个MapTask；1个ReduceTask；1个mrAppMaster 平均每个节点运行10个 / 3台 ≈ 3个任务（4    3    3） HDFS参数调优 （1）修改：hadoop-env.sh export HDFS_NAMENOD
阅读更多...
嵌入式QT开发：构建高效智能的嵌入式系统

嵌入式QT开发：构建高效智能的嵌入式系统

摘要： 本文深入探讨了嵌入式 QT 相关的各个方面。从 QT 框架的基础架构和核心概念出发，详细阐述了其在嵌入式环境中的优势与特点。文中分析了嵌入式 QT 的开发环境搭建过程，包括交叉编译工具链的配置等关键步骤。进一步探讨了嵌入式 QT 的界面设计与开发，涵盖了从基本控件的使用到复杂界面布局的构建。同时也深入研究了信号与槽机制在嵌入式系统中的应用，以及嵌入式 QT 与硬件设备的交互，包括输入输出设
阅读更多...
OpenHarmony鸿蒙开发（ Beta5.0）无感配网详解

OpenHarmony鸿蒙开发（ Beta5.0）无感配网详解

1、简介 无感配网是指在设备联网过程中无需输入热点相关账号信息，即可快速实现设备配网，是一种兼顾高效性、可靠性和安全性的配网方式。 2、配网原理 2.1 通信原理 手机和智能设备之间的信息传递，利用特有的NAN协议实现。利用手机和智能设备之间的WiFi 感知订阅、发布能力，实现了数字管家应用和设备之间的发现。在完成设备间的认证和响应后，即可发送相关配网数据。同时还支持与常规Sof
阅读更多...
活用c4d官方开发文档查询代码

活用c4d官方开发文档查询代码

当你问AI助手比如豆包，如何用python禁止掉xpresso标签时候，它会提示到 这时候要用到两个东西。https://developers.maxon.net/论坛搜索和开发文档 比如这里我就在官方找到正确的id描述 然后我就把参数标签换过来
阅读更多...
Linux_kernel驱动开发11

Linux_kernel驱动开发11

一、改回nfs方式挂载根文件系统         在产品将要上线之前，需要制作不同类型格式的根文件系统         在产品研发阶段，我们还是需要使用nfs的方式挂载根文件系统         优点：可以直接在上位机中修改文件系统内容，延长EMMC的寿命         【1】重启上位机nfs服务         sudo service nfs-kernel-server resta
阅读更多...
【区块链 + 人才服务】区块链集成开发平台 | FISCO BCOS应用案例

【区块链 + 人才服务】区块链集成开发平台 | FISCO BCOS应用案例

随着区块链技术的快速发展，越来越多的企业开始将其应用于实际业务中。然而，区块链技术的专业性使得其集成开发成为一项挑战。针对此，广东中创智慧科技有限公司基于国产开源联盟链 FISCO BCOS 推出了区块链集成开发平台。该平台基于区块链技术，提供一套全面的区块链开发工具和开发环境，支持开发者快速开发和部署区块链应用。此外，该平台还可以提供一套全面的区块链开发教程和文档，帮助开发者快速上手区块链开发。
阅读更多...
Vue3项目开发——新闻发布管理系统（六）

Vue3项目开发——新闻发布管理系统（六）

文章目录 八、首页设计开发1、页面设计2、登录访问拦截实现3、用户基本信息显示①封装用户基本信息获取接口②用户基本信息存储③用户基本信息调用④用户基本信息动态渲染 4、退出功能实现①注册点击事件②添加退出功能③数据清理 5、代码下载 八、首页设计开发 登录成功后，系统就进入了首页。接下来，也就进行首页的开发了。 1、页面设计 系统页面主要分为三部分，左侧为系统的菜单栏，右侧
阅读更多...
Node.js学习记录（二）

Node.js学习记录（二）

目录 一、express 1、初识express 2、安装express 3、创建并启动web服务器 4、监听 GET&POST 请求、响应内容给客户端 5、获取URL中携带的查询参数 6、获取URL中动态参数 7、静态资源托管 二、工具nodemon 三、express路由 1、express中路由 2、路由的匹配 3、路由模块化 4、路由模块添加前缀 四、中间件
阅读更多...
v0.dev快速开发

v0.dev快速开发

探索v0.dev：次世代开发者之利器 今之技艺日新月异，开发者之工具亦随之进步不辍。v0.dev者，新兴之开发者利器也，迅速引起众多开发者之瞩目。本文将引汝探究v0.dev之基本功能与优势，助汝速速上手，提升开发之效率。 何谓v0.dev？ v0.dev者，现代化之开发者工具也，旨在简化并加速软件开发之过程。其集多种功能于一体，助开发者高效编写、测试及部署代码。无论汝为前端开发者、后端开发者
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2