内核OpenCL中FIFO实现的最佳方法

目标：在OpenCL中实现下面显示的图表。 OpenCl内核需要的主要function是将系数数组和临时数组相乘，然后将所有这些值累加到最后。 （这可能是最耗时的操作，并行性在这里真的很有帮助）。

我正在使用内核的辅助函数进行乘法和加法（我希望这个函数也是并行的）。

图片说明：

一次一个 ，将值传递到与系数数组大小相同的数组（临时数组）中。 现在， 每次将单个值传递到此数组时，临时数组将与系数数组并行相乘，然后将每个索引的值连接成一个单独的元素。 这将继续，直到输入数组到达它的最后一个元素。

我的代码怎么了？

对于来自输入的60个元素，它需要超过8000毫秒!! 我总共需要输入120万个输入。我知道有一个更好的解决方案来做我正在尝试的事情。 这是我的代码如下。

以下是我知道他编码肯定是错误的一些事情。 当我尝试将系数值与临时数组相乘时，它会崩溃。 这是因为global_id。 我想要这条线做的只是简单地将两个数组并行。

我试图找出为什么花了这么长时间来执行FIFOfunction，所以我开始评论线路。 我首先开始评论除FIFO函数的第一个for循环之外的所有内容。 结果这耗时50毫秒。 然后，当我取消注释下一个循环时，它跳转到8000毫秒。 因此，延迟将与数据传输有关。

我可以在OpenCl中使用寄存器移位吗？ 也许对整数数组使用一些逻辑移位方法？ （我知道有一个’>>’运算符）。

float constant temp[58]; float constant tempArrayForShift[58]; float constant multipliedResult[58]; float fifo(float inputValue, float *coefficients, int sizeOfCoeff) { //take array of 58 elements (or same size as number of coefficients) //shift all elements to the right one //bring next element into index 0 from input //multiply the coefficient array with the array thats the same size of coefficients and accumilate //store into one output value of the output array //repeat till input array has reached the end int globalId = get_global_id(0); float output = 0.0f; //Shift everything down from 1 to 57 //takes about 50ms here for(int i=1; i<58; i++){ tempArrayForShift[i] = temp[i]; } //Input the new value passed from main kernel. Rest of values were shifted over so element is written at index 0. tempArrayForShift[0] = inputValue; //Takes about 8000ms with this loop included //Write values back into temp array for(int i=0; i<58; i++){ temp[i] = tempArrayForShift[i]; } //all 58 elements of the coefficient array and temp array are multiplied at the same time and stored in a new array //I am 100% sure this line is crashing the program. //multipliedResult[globalId] = coefficients[globalId] * temp[globalId]; //Sum the temp array with each other. Temp array consists of coefficients*fifo buffer for (int i = 0; i < 58; i ++) { // output = multipliedResult[i] + output; } //Returned summed value of temp array return output; } __kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) { //Initialize the temporary array values to 0 for (int i = 0; i < 58; i ++) { temp[i] = 0; tempArrayForShift[i] = 0; multipliedResult[i] = 0; } //fifo adds one element in and calls the fifo function. ALL I NEED TO DO IS SEND ONE VALUE AT A TIME HERE. for (int i = 0; i < 60; i ++) { Output[i] = fifo(Array[i], coefficients, 58); } } 

我很长一段时间都遇到过OpenCl这个问题。 我不确定如何一起实现并行和顺序指令。

我想到的另一种选择

在主cpp文件中，我正在考虑在那里实现fifo缓冲区并让内核进行乘法和加法。 但这意味着我必须在循环中将内核调用1000次以上。 这会是更好的解决方案吗？ 或者它只是完全没有效率。

Answer 1

要从GPU中获得良好的性能，您需要将您的工作并行化为多个线程。 在您的代码中，您只使用单个线程，并且每个线程的GPU速度非常慢，但如果许multithreading同时运行，则可能非常快。 在这种情况下，您可以为每个输出值使用单个线程。 您实际上不需要通过数组移动值：对于每个输出值，考虑58个值的窗口，您可以从内存中获取这些值，将它们与系数相乘并写回结果。

一个简单的实现是（使用与输出值一样多的线程启动）：

 __kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) { int globalId = get_global_id(0); float sum=0.0f; for (int i=0; i< 58; i++) { float tmp=0; if (globalId+i > 56) { tmp=Array[i+globalId-57]*coefficient[57-i]; } sum += tmp; } output[globalId]=sum; } 

这并不完美，因为它生成的内存访问模式对于GPU来说并不是最佳的。 缓存可能会有所帮助，但显然有很多优化空间，因为值会重复使用几次。 您尝试执行的操作称为卷积（1D）。 NVidia在其GPU计算SDK中有一个名为oclConvolutionSeparable的2D示例，该示例显示了优化版本。 您可以使用他们的convolutionRows内核进行一维卷积。

Answer 2

这是你可以尝试的另一个内核。 有很多同步点（障碍），但这应该表现得相当好。 65项工作组不是最优的。

步骤：

1. 将本地值初始化为0
2. 将系数复制到局部变量

循环输出元素以计算：

3. 转移现有元素（仅限工作项> 0）
4. 复制新元素（仅限工作项0）
5. 计算点积
   5A。 乘法 – 每个工作项一个
   5B。 减少循环来计算总和
6. 将点积复制到输出（仅限WI 0）
7. 最后的障碍

代码：

 __kernel void lowpass(__global float *Array, __constant float *coefficients, __global float *Output, __local float *localArray, __local float *localSums){ int globalId = get_global_id(0); int localId = get_local_id(0); int localSize = get_local_size(0); //1 init local values to 0 localArray[localId] = 0.0f //2 copy coefficients to local //don't bother with this id __constant is working for you //requires another local to be passed in: localCoeff //localCoeff[localId] = coefficients[localId]; //barrier for both steps 1 and 2 barrier(CLK_LOCAL_MEM_FENCE); float tmp; for(int i = 0; i< outputSize; i++) { //3 shift elements (+barrier) if(localId > 0){ tmp = localArray[localId -1] } barrier(CLK_LOCAL_MEM_FENCE); localArray[localId] = tmp //4 copy new element (work item 0 only, + barrier) if(localId == 0){ localArray[0] = Array[i]; } barrier(CLK_LOCAL_MEM_FENCE); //5 compute dot product //5a multiply + barrier localSums[localId] = localArray[localId] * coefficients[localId]; barrier(CLK_LOCAL_MEM_FENCE); //5b reduction loop + barrier for(int j = 1; j < localSize; j <<= 1) { int mask = (j << 1) - 1; if ((localId & mask) == 0) { localSums[local_index] += localSums[localId +j] } barrier(CLK_LOCAL_MEM_FENCE); } //6 copy dot product (WI 0 only) if(localId == 0){ Output[i] = localSums[0]; } //7 barrier //only needed if there is more code after the loop. //the barrier in #3 covers this in the case where the loop continues //barrier(CLK_LOCAL_MEM_FENCE); } } 

更多的工作组呢？
这稍微简化了一个1x65工作组计算机的整个1.2M输出。 要允许多个工作组，可以使用/ get_num_groups（0）计算每个组应该执行的工作量（workAmount），并调整i for循环：

 for (i = workAmount * get_group_id(0); i< (workAmount * (get_group_id(0)+1) -1); i++) 

步骤＃1也必须更改为初始化为localArray的正确起始状态，而不是全0。

  //1 init local values if(groupId == 0){ localArray[localId] = 0.0f }else{ localArray[localSize - localId] = Array[workAmount - localId]; } 

这两个更改应该允许您使用更多数量的工作组; 我建议设备上有多个计算单元。 尽量保持每组数量的工作量。 玩弄这个，有时在高级别上看起来最优的内核在内核运行时会对内核产生不利影响。

好处
在这个内核的几乎每一点上，工作项都有一些事情要做。 只有少于100％的项目正在工作的时间是在步骤5b中的减少循环期间。 在这里阅读更多关于为什么这是一件好事。

缺点
障碍只会通过障碍的性质来减缓内核：暂停工作项直到其他人达到这一点。 也许有一种方法可以用更少的障碍来实现这一点，但我仍然认为这是最佳的，因为你试图解决的问题。
每组没有更多工作项目的空间，65不是一个非常优化的大小。 理想情况下，你应该尝试使用2的幂，或64的倍数。但这不会是一个大问题，因为内核中有很多障碍使得它们都经常等待。