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

文章目录

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

<2022-04-27 周三>

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

不管什么事儿看来都怕琢磨，如果连做梦都能梦到你正在琢磨的事儿，估计离成功也就不远了。似乎目前已经达到了最好的效果，离目标越来越近了。

  1. 要理解clEnqueueNDRangeKernel()函数的第五第六个参数意义，但目前为止只能说暂时理解了。
cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,cl_kernel kernel,cl_uint work_dim,const size_t *global_work_offset,const size_t *global_work_size,const size_t *local_work_size,cl_uint num_events_in_wait_list,const cl_event *event_wait_list,cl_event *event)
  1. 第五第六个参数要结合__attribute__，否则无法调用kernel函数。
__kernel __attribute__((reqd_work_group_size(256, 1, 1)))
  1. 来回理解ResizeImage()和ScaleImage()函数的实现，对它们的内部逻辑流程了如指掌了可以说。
  2. 我选择以ResizeHorizontalFilter()为模板修改，之所以ResizeImage()处理速度慢，因为它的处理效果好，且有多种过滤效果可供选择，ResizeHorizontalFilter()是处理水平方向缩放，所以它将读入一整行原图像素，这正好和ScaleImage()的最内层循环处理方式相同。
  3. ResizeHorizontalFilter()的最内层循环（如下），因为有累加操作，所以这正是处理水平缩放的操作。
for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++)
{/* float weight = getResizeFilterWeight(resizeFilterCubicCoefficients,(ResizeWeightingFunctionType) resizeFilterType,(ResizeWeightingFunctionType) resizeWindowType,resizeFilterScale, resizeFilterWindowSupport,resizeFilterBlur, scale*(start + i - bisect + 0.5)); */float weight = getResizeFilterWeightForGM(resizeFilterType,scale*(start + i - bisect + 0.5), support);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);if (matte_or_cmyk != 0){cp.w = (float) *(p + 3);// float alpha = weight * QuantumScale * cp.w;// error: use of type 'double' requires cl_khr_fp64 supportfloat alpha = weight * (1 - (float) cp.w / 255);filteredPixel.x += alpha * cp.x;filteredPixel.y += alpha * cp.y;filteredPixel.z += alpha * cp.z;filteredPixel.w += weight * cp.w;gamma += alpha;}elsefilteredPixel += ((float4) weight)*cp;density += weight;
}
  1. “如何写ScaleImage()的硬件加速函数（五）”的问题在于没有办法处理图片下半部分（如何缩小一半的话），这里主要是因为y变量的限定（代码如下），因为传入kernel函数的gsize[1]=resizedRows;被限定的死死的。
const unsigned int pos = getPixelIndex(4, inputColumns, cacheRangeStartX, y);
  1. 目前只考虑按比例缩放，所以这里的y需要除以缩放比。
  2. 缩小后图片如果垂直方向相间着黑色宽竖条，那可能是因为numCachedPixels参数没有计算正确，这正可以修复“如何写ScaleImage()的硬件加速函数（五）”中的scale_ratio变量。
numCachedPixels=(int) ceil((pixelPerWorkgroup-1)/xFactor+2*support);
  1. 附上目前代码：
static MagickBooleanType scaleFilter(MagickCLDevice device,cl_command_queue queue,const Image *image,Image *filteredImage,cl_mem imageBuffer,cl_uint matte_or_cmyk,cl_uint columns,cl_uint rows,cl_mem scaledImageBuffer,cl_uint scaledColumns,cl_uint scaledRows,ExceptionInfo *exception)
{cl_kernelscaleKernel;cl_intstatus;const unsigned intworkgroupSize = 256;floatscale;intnumCachedPixels;MagickBooleanTypeoutputReady;size_tgammaAccumulatorLocalMemorySize,gsize[2],i,imageCacheLocalMemorySize,pixelAccumulatorLocalMemorySize,lsize[2],totalLocalMemorySize,weightAccumulatorLocalMemorySize;unsigned intchunkSize,pixelPerWorkgroup;scaleKernel=NULL;outputReady=MagickFalse;scale=(float) scaledColumns/columns; // TODO(ocl)if (scaledColumns < workgroupSize){chunkSize=32;pixelPerWorkgroup=32;}else{chunkSize=workgroupSize;pixelPerWorkgroup=workgroupSize;}DisableMSCWarning(4127)while(1)
RestoreMSCWarning{/* calculate the local memory size needed per workgroup */numCachedPixels=(int) ceil((pixelPerWorkgroup-1)/scale+2*(0.5+MagickEpsilon)); // TODO(ocl)imageCacheLocalMemorySize=numCachedPixels*sizeof(CLQuantum)*4;totalLocalMemorySize=imageCacheLocalMemorySize;/* local size for the pixel accumulator */pixelAccumulatorLocalMemorySize=chunkSize*sizeof(cl_float4);totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;/* local memory size for the weight accumulator */weightAccumulatorLocalMemorySize=chunkSize*sizeof(float);totalLocalMemorySize+=weightAccumulatorLocalMemorySize;/* local memory size for the gamma accumulator */gammaAccumulatorLocalMemorySize=chunkSize*sizeof(float);totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;if (totalLocalMemorySize <= device->local_memory_size)break;else{pixelPerWorkgroup=pixelPerWorkgroup/2;chunkSize=chunkSize/2;if ((pixelPerWorkgroup == 0) || (chunkSize == 0)){/* quit, fallback to CPU */goto cleanup;}}}scaleKernel=AcquireOpenCLKernel(device,"ScaleFilter");if (scaleKernel == (cl_kernel) NULL){(void) OpenCLThrowMagickException(device,exception,GetMagickModule(),ResourceLimitWarning,"AcquireOpenCLKernel failed.", ".");goto cleanup;}i=0;status =SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_mem),(void*)&imageBuffer);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_uint),(void*)&matte_or_cmyk);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_uint),(void*)&columns);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_uint),(void*)&rows);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_mem),(void*)&scaledImageBuffer);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_uint),(void*)&scaledColumns);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(cl_uint),(void*)&scaledRows);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(float),(void*)&scale);status|=SetOpenCLKernelArg(scaleKernel,i++,imageCacheLocalMemorySize,NULL);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(int),&numCachedPixels);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(unsigned int),&pixelPerWorkgroup);status|=SetOpenCLKernelArg(scaleKernel,i++,sizeof(unsigned int),&chunkSize);status|=SetOpenCLKernelArg(scaleKernel,i++,pixelAccumulatorLocalMemorySize,NULL);status|=SetOpenCLKernelArg(scaleKernel,i++,weightAccumulatorLocalMemorySize,NULL);status|=SetOpenCLKernelArg(scaleKernel,i++,gammaAccumulatorLocalMemorySize,NULL);if (status != CL_SUCCESS){(void) OpenCLThrowMagickException(device,exception,GetMagickModule(),ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");goto cleanup;}gsize[0] = (scaledColumns + pixelPerWorkgroup - 1) / pixelPerWorkgroup *workgroupSize;gsize[1] = scaledRows;lsize[0] = workgroupSize;lsize[1] = 1;outputReady=EnqueueOpenCLKernel(queue,scaleKernel,2,(const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,exception);cleanup:if (scaleKernel != (cl_kernel) NULL)ReleaseOpenCLKernel(scaleKernel);return(outputReady);
}

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){WriteAllChannels(filteredImage, 4, filteredColumns, chunkStartX + itemID, y, filteredPixel);}}
}
)

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/673136.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

c++之说_12|模板

c++之说_12|模板

关于模板&#xff0c;至少我们要先了解几个概念 一&#xff1a;函数模板 二&#xff1a;类模板 三&#xff1a;模板特化 四&#xff1a;形参参数包 模板覆盖的东西太多 我目前也不了解太多 函数模板 语法 template<typename 类型名,typename 类型名,typename ...多参…
阅读更多...
【JavaScript学习路线——详细讲解】

【JavaScript学习路线——详细讲解】

JavaScript学习路线 1. 介绍2. JavaScript基础3. HTML/CSS与JavaScript结合4. 进阶JavaScript5. 前端框架与库6. API与HTTP通信7. 移动端和PWA8. Node.js和后端开发9. 测试和部署10. 全栈项目实践11. 持续学习和社区参与 1. 介绍 学习JavaScript的路线可以分为几个主要阶段&am…
阅读更多...
【水文】计算斐波那契数列的第n项

【水文】计算斐波那契数列的第n项

#include <stdio.h> // 函数声明 int fibonacci(int n); int main() { // 输入正整数n int n; printf("请输入一个正整数 n&#xff1a;"); scanf("%d", &n); // 调用函数计算斐波那契数列的第n项并输出结果 int result …
阅读更多...
CopyOnWriteArrayList底层原理全面解析【建议收藏】

CopyOnWriteArrayList底层原理全面解析【建议收藏】

简介 CopyOnWriteArrayList是Java中的一个线程安全的集合类&#xff0c;是ArrayList线程安全版本&#xff0c;主要通过Copy-On-Write&#xff08;写时复制&#xff0c;简称COW&#xff09;机制来保证线程安全。 Copy-On-Write机制核心思想&#xff1a;向一个数组中添加数据时…
阅读更多...
LabVIEW动平衡测试与振动分析系统

LabVIEW动平衡测试与振动分析系统

LabVIEW动平衡测试与振动分析系统 介绍了利用LabVIEW软件和虚拟仪器技术开发一个动平衡测试与振动分析系统。该系统旨在提高旋转机械设备的测试精度和可靠性&#xff0c;通过精确测量和分析设备的振动数据&#xff0c;以识别和校正不平衡问题&#xff0c;从而保证机械设备的高…
阅读更多...
Leetcode 322 零钱兑换

Leetcode 322 零钱兑换

题意理解&#xff1a; 给你一个整数数组 coins &#xff0c;表示不同面额的硬币&#xff1b;以及一个整数 amount &#xff0c;表示总金额。 计算并返回可以凑成总金额所需的 最少的硬币个数 。如果没有任何一种硬币组合能组成总金额&#xff0c;返回 -1 。 你可以认为每种硬币…
阅读更多...
Springboot集成jasypt实现配置文件加密

Springboot集成jasypt实现配置文件加密

Jasypt它提供了单密钥对称加密和非对称加密两种加密方式。 单密钥对称加密&#xff1a;一个密钥加盐&#xff0c;可以同时用作内容的加密和解密依据&#xff1b; 非对称加密&#xff1a;使用公钥和私钥两个密钥&#xff0c;才可以对内容加密和解密&#xff1b; 我们以单密钥对称…
阅读更多...
前端 reduce()用法总结

前端 reduce()用法总结

定义 reduce()方法对数组中的每个元素执行一个由您提供的reducer函数(升序执行)&#xff0c;将其结果汇总为单个返回值。语法为&#xff1a; array.reduce(function(accumulator, currentValue, currentIndex, arr), initialValue); /*accumulator: 必需。累计器currentValu…
阅读更多...
5 步轻松上手，教你从 0 到 1 落地 Jmeter 接口自动化脚本！

5 步轻松上手，教你从 0 到 1 落地 Jmeter 接口自动化脚本！

Jmeter是进行接口测试的一款非常主流的工具&#xff0c;但绝大部分测试工程师&#xff0c;对于Jmeter接口测试脚本整理都是一知半解的。今天这篇文章&#xff0c;就以一个金融项目中接口为例&#xff0c;通过简单5步&#xff0c;教大家如何0代码编写Jmeter接口自动化脚本&#…
阅读更多...
CPU和GPU有什么区别，玩游戏哪个更重要？

CPU和GPU有什么区别，玩游戏哪个更重要？

大家好&#xff01;今天我们要聊的话题是CPU和GPU&#xff0c;它们在电脑中扮演着重要的角色&#xff0c;虽然看起来只是两个简单的缩写&#xff0c;但它们的功能和影响是截然不同的&#xff01; 那么&#xff0c;究竟CPU和GPU有什么区别呢&#xff1f;在玩游戏时&#xff0c;…
阅读更多...
Linux 系统开启网络服务

Linux 系统开启网络服务

首先&#xff0c;大家新装的linux系统可能都没有安装vim工具&#xff0c;所以打开文件的方式是 vi /etc/sysconfig/network-scripts/ifcfg-ens33在这个界面把onboot改为yes&#xff0c;我这里是设置完的。然后通过下面语句重新启动服务就可以了。 service network restartcen…
阅读更多...
2024.2.7日总结（小程序开发4）

2024.2.7日总结（小程序开发4）

页面导航 页面导航是页面之间的相互跳转&#xff1a; <a>链接location.href 小程序中实现页面导航的两种方式&#xff1a; 声明式导航 在页面上声明一个<navigator>导航组件通过点击<navigator>组件实现页面跳转 编程式导航 调用小程序的导航API&…
阅读更多...
飞天使-k8s知识点15-kubernetes散装知识点4-CNI网络插件与kubectl

飞天使-k8s知识点15-kubernetes散装知识点4-CNI网络插件与kubectl

文章目录 CNI 网络插件安装任意节点运行kubectlAPI的版本区别与废弃API查询 CNI 网络插件安装 这里将以 Calico 为例&#xff0c;提供在 Kubernetes 1.20.6 版本上安装 CNI 插件的步骤。请注意&#xff0c;具体的步骤可能会因 CNI 插件的类型和你的特定环境而略有不同。设置 Ku…
阅读更多...
BaseMapper中提供的方法(17种CRUD)

BaseMapper中提供的方法(17种CRUD)

BaseMapper封装的17种增删改查方法 MybatisPlus框架中mapper层继承了BaseMapper接口&#xff0c;该接口中封装了常用的增删改查方法&#xff0c;共有17种&#xff0c;以下是方法的详情介绍 首先需要明确的括号内的一些对象定义 泛型T&#xff1a;实体类类型Param注解&#x…
阅读更多...
JavaScript 设计模式之单例模式

JavaScript 设计模式之单例模式

单例模式 常规单例 单例模式我们在日常使用中还是非常多的&#xff0c;比如常见的 jQuery&#xff0c;prototype&#xff0c;vue等都是属于单例模式&#xff0c;我们在使用 new Vue 的时候&#xff0c;返回的也会是同一个实例的&#xff0c;简单实现 // 方式一 let Car func…
阅读更多...
mybatis-plus循环处理多个条件的 or 查询

mybatis-plus循环处理多个条件的 or 查询

我们一般用 mybatis-plus 的提供的 api 接口处理 List、Set 作为条件查询的时候&#xff0c;都会使用 in&#xff0c;例如 &#xff08;Student 类省略 没啥好些的&#xff09;&#xff1a; LambdaQueryWrapper<Student> queryWrapper new QueryWrapper<Student>…
阅读更多...
iPhone解锁 AnyMP4 iPhone Unlocker

iPhone解锁 AnyMP4 iPhone Unlocker

AnyMP4 iPhone Unlocker是一款功能强大的iPhone解锁软件&#xff0c;旨在帮助用户轻松解决iPhone密码忘记、设备锁定等问题。无论是屏幕密码、指纹解锁还是Face ID&#xff0c;该软件都能提供有效的解决方案。 这款软件支持多种iPhone型号&#xff0c;包括最新的iPhone 14系列…
阅读更多...
2.3_9 吸烟者问题

2.3_9 吸烟者问题

2.3_9 吸烟者问题 问题描述 问题分析 假设一个系统有三个抽烟者进程和一个供应者进程。每个抽烟者不停地卷烟并抽掉它&#xff0c;但是要卷起并抽掉一支烟&#xff0c;抽烟者需要有三种材料&#xff1a;烟草、纸和胶水。三个抽烟者中&#xff0c;第一个拥有烟草、第二个拥有纸…
阅读更多...
泛娱乐社交出海洞察，Flat Ads解锁海外增长新思路

泛娱乐社交出海洞察，Flat Ads解锁海外增长新思路

摘要:解读泛娱乐社交应用出海现状与趋势,解锁“掘金”泛娱乐社交出海赛道新思路。 根据全球舆情监测机构 Meltwater 和社交媒体机构We are Social最新发布数据显示,全球社交媒体活跃用户数量已突破50亿,约占世界人口总数62.5%。庞大的用户数量意味着广阔的增量空间,目前,随着全…
阅读更多...
使用HCPpipelines分割皮层

使用HCPpipelines分割皮层

前段时间阅读了一篇文献，文章的做法我比较感兴趣，所以打算学习一下文献的做法。文章的最开始一部分是使用HCPpipelines对T1和T2像进行皮层分割，调用的是freesurfer6。https://github.com/Washington-University/HCPpipelines 一、工作环境准备 1.安装好FSL，版本在6.0.2以上…
阅读更多...
最新文章

Copyright @ 2022~2023