OpenCL 内核执行速度非常慢，我做错了什么吗？

问题描述 投票：0回答：1

我为 Open CL 编写了一个简单移动平均滤波器，它的执行速度比在 CPU 上慢 2 到 3 个数量级。

简单移动平均滤波器是一个串行操作，我在调用 

cl::NDRange(1)
时使用 
enqueueNDRangeKernel()
执行。

我希望在 GPU 而不是 CPU 上执行此操作的原因是因为在该内核之前有一些大型数组的并行处理，效果很好，并且在移动平均滤波器之后我只需要大约 1 in 100 个样本，这样我可以显着减少从 GPU 复制回主机的数据量，从而加快操作速度。

我希望 GPU 不一定会击败 CPU，但不会比 CPU 慢几倍，我会通过减少数据传输来弥补这一点。

我的内核：

__kernel void SingleMovingAverage(__global double* InputSamples, 
                                  __global double* FilteredSamples, 
                  __global double* CircularBuffer,
                  double RunningTotal
                 )
{
    int CircularBufferPointerCtr = 0;   
    
    for(int sampleCtr = 0; sampleCtr < 262144; sampleCtr++)
    {
        RunningTotal += InputSamples[sampleCtr];
        RunningTotal -= CircularBuffer[CircularBufferPointerCtr];
        FilteredSamples[sampleCtr] = RunningTotal;
        
        CircularBuffer[CircularBufferPointerCtr] = InputSamples[sampleCtr];
        
        CircularBufferPointerCtr++;
        CircularBufferPointerCtr &= 0xFF;
    }
}

也许 GPU 在执行单个串行任务时会慢很多？

opencl
1个回答
0
投票

除了令人尴尬的并行（以方便的向量格式对数据进行大量相同的、无依赖性的操作）之外，您总是需要期望对算法进行一些修改，以便有任何机会使其在GPU。

当然，移动平均算法的简单实现是严格串行的，但这并不意味着不可能并行。在优化串行代码的世界中，您需要放弃的一件事是，“增加”总操作计数以“减少”顺序相关操作的数量是完全正常的。另一件事是代码复杂性通常会上升。 看起来您正在获取 262144 (N) 个输入样本 (

I

) 并生成 262144 (

N
) 个滤波输出 (
F
)，其中每个输出是最后 256 个窗口的总和(
M
) 采样直至当前输入（包括当前输入）。
我突然想到的一个并行化的想法是使用分块，并通过多次处理来解决问题。本质上，我们将进行以下观察：
我们在 
F[j] = sum(I[j-255..j])
、
F[255] = sum(I[0..255])
、
F[256] = sum(I[1..256])
 等之后。（具有适当的边界条件，使得 
I
 的负索引产生零。）这可以重写为：


T[j] = sum(I[j*32..j*32+31])
（
    T[0] = sum(I[0..31])
  • 、
    T[1] = sum(I[32..63])
    等
    ）

    F[255] = sum(T[0..7]) = T[0] + sum(T[1..7])
     或
  • F[255] = sum(I[0..31]) + sum(T[1..7])
    

    F[256] = sum(I[1..31]) + sum(T[1..7]) + sum(I[0..0])
    • 

  • F[257] = sum(I[2..31]) + sum(T[1..7]) + sum(I[0..1])
    • 

  • 等等
    

    • 我们可以重复使用这些
  • sum(T[1..7])
    • 表达式！前面的位（
sum(I[2..31])
等）是后缀和，而后面的位（
sum(I[0..0])
）是前缀和。


选择块大小 (
C
)，它是窗口大小的适当较小的 pow2 因子。使用问题大小的平方根 (sqrt(M) = 16) 通常很容易实现，并且通常在理论上是最佳的，但在这种情况下可能会导致您难以优化或修改窗口大小，并且在实践中也不是最佳的；但对于概念证明来说它可能已经足够好了。我将在下面的示例中选择 
    C=16
  1. ：
    分配临时缓冲区 
    T
     来保存 
    2. N/C
  2.  和。您将在内核之间传递它。
    提交第一个内核，其中包含 
    N
     工作项和组大小 
    3. C
  3. 。这对 16 个相邻输入样本的组执行并行归约和。 
    InputSamples[0..C-1] -> T[0]
    、
    InputSamples[C..2*C-1]
     -> 
    T[1]
     等
    再次提交第二个内核，其中包含 
    N
     工作项和 
    4. C
  4.  的组大小。

    首先，对线程组生成的 
    all
     输出共享的 
      M/C-1
    •  块执行并行求和，我们将结果称为 
      T
      。
      在第一个共享 Ts 之前立即对 
      S
       输入样本执行并行后缀和缩减 
      • C
    • 。对最后共享的 
      T
       之后的 
      P
       输入值执行并行 
      prefix
       总和缩减 C。
      每个过滤后的输出值为 
      T
      。
      

      • 

    • 

      串行算法使用线性缩放的操作数。 (O(
      S[i] + Ts + P[i]
      • ))
    

    5. 并行算法的计算复杂度大约为 
N
，但常数因子较小，并且扩展到 
O(N * (3*log(C) + log(M/C))
 个执行单元，
N
 操作和因此时间复杂度约为 
O(3*log(C) + log(M/C))
。您可能不会有 262144 个核心，实际上整个过程不会那么高效，但您仍然应该看到性能的大幅跃升。

该算法的一个缺点是我们在两个不同的内核中读取整个输入数组，这意味着使用更多的内存带宽。分析和性能计数器会告诉您这是否足以导致目标硬件上的 ALU 资源不足。 （如果您的数据大小确实只有 2MB，那么无论如何它都应该适合缓存。）
    



      

      
    

  

  

    
最新问题


  





  

    © www.soinside.com 2019 - 2024. All rights reserved.