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

Question

我一直在使用 C# + Cudafy(C# -> CUDA 或 OpenCL 转换器)开发的波形模拟器运行良好，除了运行 OpenCL CPU 版本(Intel 驱动程序，15) "MacBook Pro Retina i7 2.7GHz，GeForce 650M(开普勒，384 核))大约是 GPU 版本的四倍。

(无论我使用 CL 还是 CUDA GPU 后端都会发生这种情况。OpenCL GPU 和 CUDA 版本的性能几乎相同。)

澄清一下，对于示例问题:

• 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 是错误的，但即使由此产生的流水线停顿也不应导致如此极端的性能增量。

唯一让我吃惊的是，我使用的是一个蹩脚的网格/ block 分配方案——即，网格是要更新的数组的大小，每个 block 是一个线程。我确信这对性能有一些影响，但我看不到它导致它成为在 CPU 上运行的 CL 代码速度的 1/4。啊!

Answer 1

回答此问题以将其从未回答列表中删除。

发布的代码表明内核启动正在指定 1 个(事件)线程的线程 block 。这不是编写快速 GPU 代码的方法，因为它会使大部分 GPU 功能闲置。

典型的线程 block 大小应至少为每个 block 128 个线程，通常越高越好，为 32 的倍数，最多为每个 block 512 或 1024 个线程，具体取决于 GPU。

GPU“喜欢”通过大量“可用”的并行工作来隐藏延迟。为每个 block 指定更多线程有助于实现这一目标。 (在网格中拥有相当多的线程 block 也可能有所帮助。)

此外，GPU 以 32 个为一组执行线程。指定每个 block 仅 1 个线程或 32 的非倍数将在每个执行的线程 block 中留下一些空闲执行槽。每个 block 1 个线程特别糟糕。