Cuda-OpenCL CPU比OpenCL或CUDA GPU版本快4倍

问题描述：

我一直在用C＃+ Cudafy（C＃ - > CUDA或OpenCL translator）工作的波形模拟器效果很好，除了运行OpenCL CPU版本（Intel驱动，15“MacBook Pro Retina i7 2.7GHz，GeForce 650M（开普勒，384内核））大约是GPU版本的四倍。。后端OpenCL的GPU和CUDA版本执行几乎相同）Cuda-OpenCL CPU比OpenCL或CUDA GPU版本快4倍

为了澄清，用于样品的问题：

OpenCL的CPU 1200赫兹
的OpenCL GPU 320赫兹
CUDA GPU - 〜330赫兹

我茫然地解释为什么CPU版本将更快比GPU。在这种情况下，CPU和GPU上执行的内核代码（在CL情况下）是相同的。我在初始化期间选择CPU或GPU设备，但除此之外，所有内容都是相同的。

编辑

下面是启动内核之一的C＃代码。（其他是非常相似的。）

public override void UpdateEz(Source source, float Time, float ca, float cb) 
    { 
     var blockSize = new dim3(1); 
     var gridSize = new dim3(_gpuEz.Field.GetLength(0),_gpuEz.Field.GetLength(1)); 

     Gpu.Launch(gridSize, blockSize) 
      .CudaUpdateEz(
       Time 
       , ca 
       , cb 
       , source.Position.X 
       , source.Position.Y 
       , source.Value 
       , _gpuHx.Field 
       , _gpuHy.Field 
       , _gpuEz.Field 
      ); 

    }

而且，这里是由Cudafy产生的相关CUDA内核函数：

extern "C" __global__ void CudaUpdateEz(float time, float ca, float cb, int sourceX, int sourceY, float sourceValue, float* hx, int hxLen0, int hxLen1, float* hy, int hyLen0, int hyLen1, float* ez, int ezLen0, int ezLen1) 
{ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    if (x > 0 && x < ezLen0 - 1 && y > 0 && y < ezLen1 - 1) 
    { 
     ez[(x) * ezLen1 + (y)] = ca * ez[(x) * ezLen1 + (y)] + cb * (hy[(x) * hyLen1 + (y)] - hy[(x - 1) * hyLen1 + (y)]) - cb * (hx[(x) * hxLen1 + (y)] - hx[(x) * hxLen1 + (y - 1)]); 
    } 
    if (x == sourceX && y == sourceY) 
    { 
     ez[(x) * ezLen1 + (y)] += sourceValue; 
    } 
}

只是为了完整性，这里是用来生成CUDA的C＃：

[Cudafy] 
    public static void CudaUpdateEz(
     GThread thread 
     , float time 
     , float ca 
     , float cb 
     , int sourceX 
     , int sourceY 
     , float sourceValue 
     , float[,] hx 
     , float[,] hy 
     , float[,] ez 
     ) 
    { 
     var i = thread.blockIdx.x; 
     var j = thread.blockIdx.y; 

     if (i > 0 && i < ez.GetLength(0) - 1 && j > 0 && j < ez.GetLength(1) - 1) 
      ez[i, j] = 
       ca * ez[i, j] 
       + 
       cb * (hy[i, j] - hy[i - 1, j]) 
       - 
       cb * (hx[i, j] - hx[i, j - 1]) 
       ; 

     if (i == sourceX && j == sourceY) 
      ez[i, j] += sourceValue; 
    }

显然，在这种内核的if是坏的，但即使所产生的流水线停顿不会引起这种极端的性能三角洲。

跳到我身上的唯一另外一件事是我使用了一个不成熟的网格/块分配方案 - 即网格是要更新的数组的大小，并且每个块都是一个线程。我确信这会对性能产生一些影响，但我看不到它导致它在CPU上运行的CL代码速度的四分之一。哎呀！

你有一些你可以共享的代码示例吗？ – 2013-05-07 23:13:53

@EricBainville当然 - 你想要C＃，CUDA或CL内核，还是什么？（这是一个半中型应用程序，我不想将20k行代码粘贴到SO中） – 2013-05-07 23:16:03

+10

我没有看到任何迹象表明cuda内核每块使用多于1个线程（没有使用'threadIdx.x'或'threadIdx.y'）。此外，启动每块指定1个线程。这意味着大约97％的GPU功能未被使用。我对cudafy了解不多，所以我不知道你是否可以控制这个，但我并不觉得cuda代码的运行速度并不快。 – 2013-05-08 00:08:10

答

回答这个得到它的解答列表。

的代码发布指示内核启动被指定1（活性）螺纹的threadblock。这不是编写快速GPU代码的方式，因为它会让大部分GPU功能闲置。

典型threadblock尺寸应为每块的至少128个线程，以及更高的是往往更好，以32的倍数，高达512或1024每块的限制，这取决于GPU。

GPU的“喜欢”具有“可用”大量的并行工作来隐藏延迟。为每个块指定更多线程有助于实现此目标。（具有在网格中的相当大的数目的threadblocks也可能有帮助。）

此外，该GPU中的32组执行线程。每个块只指定一个线程或32的非倍数会在执行的每个线程块中留下一些空闲执行槽。每块1个线程特别糟糕。

Cuda-OpenCL CPU比OpenCL或CUDA GPU版本快4倍

相关推荐