AMD OpenCL例子阅读笔记系列之BinarySearch
来源:互联网 发布:大连知达到曲阜东高铁 编辑:程序博客网 时间:2024/05/11 21:17
首先为了更好地理解,先把附带的文档大概翻一下。这个例子的主要思想是将大数据通过类似缩减树的形式对搜索范围进行缩减至最终生成结果。
介绍
实现细节
这时,只有线程2输出,同时第三个元素的值为1,意味着搜索结束。后面还讲到了如果搜索43,那么还将继续规约,每一个线程将会处理一个数据。
例如42,44,46……但是43与此都不相同,而下一次规约的大小小于1,所以可以下结论得出找不到数据。
BinarySearch_kernels.cl中有三个实现:
1)binarySearch:本文描述的算法
2)binarySearch_mulkeys:查找多个值
3) binarySearch_mulkeysConcurrent:并行查找多个值
首先是binarySearch内核:
__kernel voidbinarySearch( __global uint4 * outputArray, __const __global uint2 * sortedArray, const unsigned int findMe){ unsigned int tid = get_global_id(0); /* Then we find the elements for this thread */ uint2 element = sortedArray[tid]; /* If the element to be found does not lie between them, then nothing left to do in this thread */ if( (element.x > findMe) || (element.y < findMe)) { return; } else { /* However, if the element does lie between the lower and upper bounds of this thread's searchspace * we need to narrow down the search further in this search space */ /* The search space for this thread is marked in the output as being the total search space for the next pass */ outputArray[0].x = tid; outputArray[0].w = 1; }}输入量sortedArray为分段后的上下界信息,findMe当然就是要找的值了。输出量outputArray,利用线程号进行标记。该代码比较简单,就是在查找到的时候对输出所需信息进行记录,否则直接返回。
我们看下主机端代码:
首先是在setupBinarySearch中对测试数据进行初始化,关键代码为:
input[0] = 0; for(cl_uint i = 1; i < length; i++) { input[i] = input[i - 1] + (cl_uint) ((max * rand()) / (float)RAND_MAX); }通过随机数进行初始化,而内核的执行主要在runCLKernels函数中。函数首先通过跟所选设备的GroupSize进行比较,然后确定最终的分割数目subdivSize而对于输入量上下界的代码如下:
cl_uint *in=NULL; // Set input data status = mapBuffer( inputBuffer, in, inlength, CL_MAP_WRITE ); CHECK_ERROR(status, SDK_SUCCESS, "Failed to map device buffer.(inputBuffer)"); for(cl_uint i=0 ; i<numSubdivisions; i++) { int indexa = i*subdivSize; int indexb = (i+1)*subdivSize-1; in[2*i] = input[indexa]; in[2*i+1] = input[indexb]; } status = unmapBuffer( inputBuffer, in); CHECK_ERROR(status, SDK_SUCCESS, "Failed to unmap device buffer.(inputBuffer)");这里使用了OpenCL的内存Map技术,这种方式可以将设备内存映射到主机上得到主机端的指针,通过指针对该块内存区域进行操作。然后看一下内核参数的传递:
status = clSetKernelArg( kernel, 0, sizeof(cl_mem), (void *)&outputBuffer); CHECK_OPENCL_ERROR(status, "clSetKernelArg 0(OutputBuffer) failed.");
/* * Second argument is input buffer */ status = clSetKernelArg( kernel, 1, sizeof(cl_mem), (void *)&inputBuffer); CHECK_OPENCL_ERROR(status, "clSetKernelArg 1(inputBuffer) failed."); /* * Third is the element we are looking for */ status = clSetKernelArg( kernel, 2, sizeof(cl_uint), (void *)&findMe); CHECK_OPENCL_ERROR(status, "clSetKernelArg 2(findMe) failed.");
之所以特意把这两部分分开是因为第一个参数实际上是公用的输出参数,只需要传入一次,而第二个和为每次需要变化的,而且我认为第3个参数也应该放到循环外面好些。设置完成后就要讲该内核入队进行处理了:
cl_event ndrEvt; status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalThreads, localThreads, 0, NULL, &ndrEvt);在每次处理之后当然最关键的就是搜索范围的缩减,缩减的过程如下:
globalLowerBound = output[0]*subdivSize; globalUpperBound = globalLowerBound+subdivSize-1; subdivSize = (globalUpperBound - globalLowerBound + 1)/numSubdivisions;
output[0]从内核可以看出是每次搜索在范围内的线程进行标记,所以很容易地可以由他转换成下一次线程的上下界和分割量。而整体的运行过程可以用下面的伪代码说明一下:
while(subdivSize>1 && output[3]!=0)
{
设定新的上下界值
更新内核参数
重设output[3]=0
执行内核
读取输出值
更新搜索上下界及分割大小
}
所作这些的最终目的当然是找到我们需要的值,实际上上面也只是最终确定了一个最小比较范围,但是这样对于查找来说也已经简化了不少。具体查找如下:
for(cl_uint i=globalLowerBound; i<= globalUpperBound; i++) { if(input[i] == findMe) { elementIndex = i; globalLowerBound = i; globalUpperBound = i+1; isElementFound = 1; break; } }但是唯一不理解的时对output的Map操作怪怪的,好像应该在while循环里面进行unmap才对。
最后我们再来看下另外两个内核:
__kernel voidbinarySearch_mulkeys(__global int *keys, __global uint *input, const unsigned int numKeys, __global int *output){ int gid = get_global_id(0); int lBound = gid * 256; int uBound = lBound + 255; //int uBound for(int i = 0; i < numKeys; i++) { if(keys[i] >= input[lBound] && keys[i] <=input[uBound]) output[i]=lBound; } }
有了上面的说明,这个例子就好理解多了,只是把查找的值变成了多个,输出对应也变成多个而已。再看看并行版本的:
__kernel voidbinarySearch_mulkeysConcurrent(__global uint *keys, __global uint *input, const unsigned int inputSize, // number of input elements const unsigned int numSubdivisions, __global int *output){ int lBound = (get_global_id(0) % numSubdivisions) * (inputSize / numSubdivisions); int uBound = lBound + inputSize / numSubdivisions; int myKey = keys[get_global_id(0) / numSubdivisions]; int mid; while(uBound >= lBound) { mid = (lBound + uBound) / 2; if(input[mid] == myKey) { output[get_global_id(0) / numSubdivisions] = mid; return; } else if(input[mid] > myKey) uBound = mid - 1; else lBound = mid + 1; }}似乎复杂了不少,我们总体一看,那个while就是二分查找嘛,那上面那些参数是做什么用的?我们一个个来看。
lBound顾名思义肯定是下界,inputSize/numSubdivisions,这个是分割后的段大小,get_global_id(0)%numSubdivisions得到的是0~numSubdivisions的数。我们合起来看就是0~numSubdivision的数乘以分割后段大小,那么很显然就是分给后一段的下界。
uBound这个简单只是在lBound的基础上加上分割后的段长度。myKey是在要找的keys中取得该段数据应当查找哪个关键字。
说了这么多到底怎么个并行法,实际上就是例如:0~100,100~200,200~300三个范围的自然数,分别要找的数分别为10,120,256。那么将numSubdivisions设为3,第一个线程可以计算出计算的范围为0~100,取得要找的值为10,第二个线程要找的值为120……。这样三段线程分别可以并行查找该段内对应的key值。不知道理解的对不对。如果不对请大神看到在下面回复下。
- AMD OpenCL例子阅读笔记系列之BinarySearch
- AMD OpenCL例子阅读笔记系列之AtomicCounters
- AMD OpenCL例子阅读笔记系列之DeviceFission
- AMD OpenCL例子阅读笔记系列之Radix_Sort(一)
- AMD OpenCL例子阅读笔记系列之Radix_Sort(二)
- AMD OpenCL例子阅读笔记系列之Radix_Sort(三)
- AMD OpenCL例子阅读笔记系列之Radix_Sort(四)
- AMD OpenCL例子阅读笔记系列之Radix_Sort(五)
- AMD OpenCL例子阅读笔记系列之Radix_Sort(六)
- OpenCL 1.0 Specification阅读笔记(1)
- OpenCL 1.0 Specification阅读笔记(2)
- OpenCL 1.0 Specification阅读笔记(3)
- OpenCL 1.0 Specification阅读笔记(4)
- OpenCL 1.0 Specification阅读笔记(5)
- OpenCL 1.0 Specification阅读笔记(6)
- OpenCL 1.0 Specification阅读笔记(7)
- OpenCL 1.0 Specification阅读笔记(8)
- OpenCL 1.0 Specification阅读笔记(9)
- 二叉树的一些操作整理
- poj The Unique MST
- 盛大android面试归来
- poj 1265 Area(pick公式)
- python PIL 检查图片的完整性
- AMD OpenCL例子阅读笔记系列之BinarySearch
- word在尾注后添加致谢
- MySQL游标和控制结构代码学习
- SVN修改端口号
- JavaScript中 逗号运算符以及eval
- SSL/TLS协议运行机制的概述
- 改善java程序的151个建议 读后感一
- 将想法变为软件的过程
- 关于TreeView绑定数据库的表数据如何,代码在前两篇博文中~~