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

Question

目标:在 OpenCL 中实现下图所示。 OpenCl 内核所需的主要内容是将系数数组和临时数组相乘，然后最后将所有这些值累加为 1。 (这可能是最耗时的操作，并行性在这里非常有帮助)。

我正在为内核使用一个辅助函数来执行乘法和加法(我希望这个函数也将是并行的)。

图片描述:

一次一个，这些值被传递到与系数数组大小相同的数组(临时数组)中。现在，每次将单个值传递到该数组时，临时数组都会与系数数组并行相乘，然后将每个索引的值连接到一个元素中。这将继续，直到输入数组到达其最终元素。

我的代码会发生什么情况？

对于输入中的 60 个元素，需要超过 8000 毫秒!我总共有 120 万个输入需要传入。我知道有一个更好的解决方案可以完成我正在尝试的事情。下面是我的代码。

以下是我所知道的代码中肯定存在的一些问题。当我尝试将系数值与临时数组相乘时，它崩溃了。这是因为 global_id。我想要这行代码做的只是将两个数组并行相乘。

我试图找出为什么执行 FIFO 函数需要这么长时间，所以我开始注释行。我首先对除 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 每个线程的速度非常慢，但如果同时运行多个线程，则速度可能会非常快。在这种情况下，您可以对每个输出值使用单个线程。您实际上不需要通过数组移动值:对于每个输出值，都会考虑包含 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 内核进行 1D 卷积。