OpenCL维度跨越

上面讲到怎么在一维问题里面进行操作那么下面来讲述一下二维矩阵或者图像。在OpenCL里面这种跨越核心在于坐标系的变化。坐标系的升级在一维里我们只用到了get_global_id(0)。但是在二维中GPU会给每个工人分配(x,y)get_global_id(0),get_global_id(1)分别代表了列和行内核代码的变化假设我们要处理一张图片时:__kernel void matrix_example(__global float*data,int width) { int xget_global_id(0); int yget_global_id(1); //将二维坐标映射回一维内存地址 int indexy*widthx; //把每个像素亮度翻倍 data[index]data[index]*2.0f; }来理解一下int indexy*widthx。首先y是行号x是列号其实举例子就很好理解。假设我是第三行第二个人那么我的座号就是前面两行人总数2。这里的宽度就是每行人有多少个人。在图片中图片的宽度就是指图片的每一行有多少个像素。那如果我们处理的是数学矩阵而不是图片的话比如处理一个500*500的矩阵那么宽度就是500。因为GPU不认识行和列在它眼里我们的二维就是它的一维它只认识第几个。那么再来讲点图形在计算机里一张图片其实是由无数个颜色数值组成的如果我们考虑最简单的灰阶图片黑白照那么0代表纯黑色255OpenCL常用1.0代表纯白色中间的数值就是不同程度的灰色。那么亮度在数学上就是这个像素点的数值大小数值越大这个点看起来就越亮其实就是在调整该像素的灰度值。假设说一个像素的值是0.3那么执行了data[index]*2.0f数值翻倍视觉上就变亮了。那么整个代码就是派出很多个工人执行这段代码读取自己负责的像素乘以2后写回原位。C端的操作在C调用时最关键的变化在于clEnqueueNDRangeKernel的参数。以前我们传1维现在我们要告诉GPU这是2维。size_t global_work_size[2]{width,height};//告诉GPU我们的计算网格有多少行多少列比如说一张图片是1920*1080那么这里的width1920height1080。 clEnqueueNDRangeKernel( queue, kernel, 2, //该参数指的是维度图像是二维的 NULL, global_work_size, NULL,0,NULL,NULL );思考一个问题如果一个像素本来就很亮数值是0.8当把它乘以2后超过了最大值1.0。那么怎么在代码里面防止这种“爆表”情况按照C的想法简单的做法就是采用if语句但实际上在GPU编程里面频繁使用分支if/else可能会让工人待工。一堆工人里面一半要用if语句另一半不用的人可能就得在原地打转等另一半人干完。这被称为分支散布。为了让代码运行更平滑OpenCL提供了一些内置数学函数我们可以使用clamp或者min函数。有两种方案方案一data[index]min(data[index]*2.0f,1.0f)。方案二data[index]clamp(data[index]*2.0f,0.0f,1.0f)这个就是设定死范围了最小值和最大值之间。现在来学习一下GPU的“补齐”现在处理一个100*50的矩阵100是width50是heightglobal_size[0]对应的是width列/横向。假设总共需要总人数512*512那么GPU会把它们分成一个个小车间比如说16*16。size_t global_size[2]{512,512};//需要的工人总数 size_t local_size[2]{16,16};//每个车间的规模 clEnqueueNDRangeKernel( queue, kernel, 2, NULL, global_size, local_size,//这里如果是NULL就是GPU自动选择一个车间要有多少人 0,NULL,NULL );那如果我们的图片宽度不是16的倍数比如宽度是500但是仍要求每个车间必须是16*16的话。这个时候global_size就设置成512*512。为什么不是派500*500个工人因为500无法被16整除那么就意味着边缘的车间缺人GPU可能会因为工作组不完整而报错或拒绝执行。所以最简单的办法就是把global_work_size向上取整到16的整数倍。但是这样带来新的问题越界。假设工厂里有512*512个工人现它们去读取data[index]的任务那么多出来的人就会发生非法访问内存。那么为了避免越界问题就用if语句来加锁工人的(x,y)坐标都要小于500不是小于等于500索引从0开始__kernel void image_brightness(__global float*data,int width,int height,float brightFactor) { int xget_global_id(0); int yget_global_id(1); //加锁 if(xwidthyheight) { int indexy*widthx; data[index]min(data[index]*brightFactor,1.0f); } }之前说到分支散布会使程序效率变低但是我们还是优先考虑代码的正确性。怎么避开if有两个方案方案一数据对齐在创建图片内存的时候直接申请512*512的空间。虽然图片内容是500*500。方案二使用OpenCL的image2d_t对象而不是普通的__global float*。现在来学习跨距现在有一个问题假设我们为了对齐申请了一个宽度为512的数组来存宽度为500的图片有一个工人坐标是x0y1也就是第二行的第一个像素那么回答下面几个问题算出该人的index假设该人的index是500它读到的数据还是第二行第一个像素吗那么第一题答案是512512*10.第二题的答案是不是它读到的是第一行后面的空格而已因为第一行有512个格子其中0-499个格子才有像素值。所以在工业级开发中这个用于计算的宽度512被称作pitch或者stride跨距而图片的实际宽度500被称作width。现在有类似的问题进行巩固假设我们要处理一张1000*1000的灰度图为了GPU效率我们申请了1024个像素的内存空间。传入哪些参数__global float*data,int width,int height,int stride,float brightFactorif里面判断谁if(xwidthyheight)计算index的公式int indexy*stridex复杂一点的跨越刚刚理解了怎么处理不规则的二维数组现在我们试着理解一下矩形转置。假设我们有一张图片要把它的行和列交换。输入xy要把处理后的数据存入新的数组output那么output里面的index_out公式该怎么写首先在读取时index_iny*stride_inx。这里的stride_in是原本的宽度跨距。在写入阶段就是每个工人读取完数据对数据进行加工处理后将处理后的数据存入outputindex_outx*stride_outy。这里自己试着画图很好理解这个式子的。注意这里的stride_out对应原图的高度方向。学习三维世界现在理解了二维让我们学习怎么去处理一栋大楼get_global_id(0):房间号xget_global_id(1):第几排yget_global_id(2):楼层z在这种情况下如何去定位一个具体房间。假设大楼的宽度是W每一排有W个房间每一层楼的面积一层的总房间数是W*H。那么我们想找到第z层第y排第x个房间就要先跳过z*W*H)(y*W)x。那么在实际的GPU优化中为了内存对齐每一排有一个stride每一层也有一个slicestride。意思就是实际有用到的宽度是width实际有用到的二维高度是height但是每一排占用的空间是stride每一层占用的空间是slice_stride一整层的总空间所以最精准的定位公式就是indexz*slice_stridey*stridex。