使用 float4、opencl 时的加速答案

【问题标题】：speedup when using float4, opencl使用 float4、opencl 时的加速
【发布时间】：2016-09-26 11:27:34
【问题描述】：

我有以下 opencl 核函数来获取图像的列总和。

__kernel void columnSum(__global float* src,__global float* dst,int srcCols,
                            int srcRows,int srcStep,int dstStep)   
{

    const int x = get_global_id(0);
    srcStep >>= 2;
    dstStep >>= 2;

    if (x < srcCols)
    {
        int srcIdx = x ;
        int dstIdx = x ;

        float sum = 0;

        for (int y = 0; y < srcRows; ++y)
        {
            sum += src[srcIdx];
            dst[dstIdx] = sum;
            srcIdx += srcStep;
            dstIdx += dstStep;
        }
    }
}

这里我指定每个线程处理一列，这样很多线程可以并行获取每一列的column_sum。

我还用float4重写了上面的内核，这样每个线程可以从源图像中一次读取一行4个元素，如下图。

__kernel void columnSum(__global float* src,__global float* dst,int srcCols,
                            int srcRows,int srcStep,int dstStep)
{

    const int x = get_global_id(0);

    srcStep >>= 2;
    dstStep >>= 2;
    if (x < srcCols/4)
    {
        int srcIdx = x ;
        int dstIdx = x ;

        float4 sum = (float4)(0.0f, 0.0f, 0.0f, 0.0f);

        for (int y = 0; y < srcRows; ++y)
        {
            float4 temp2;
            temp2 = vload4(0, &src[4 * srcIdx]);
            sum = sum + temp2;

            vstore4(sum, 0, &dst[4 * dstIdx]);

            srcIdx += (srcStep/4);
            dstIdx += (dstStep/4);
        }
    }
}

在这种情况下，理论上，我认为第二个内核处理图像所消耗的时间应该是第一个内核函数所消耗时间的1/4。但是，无论图像有多大，两个内核几乎消耗相同的时间。我不知道为什么。你们能给我一些想法吗？ T

【问题讨论】：

那么瓶颈不在于信息共享部分。除了你没有得到 1/4 的加速之外，性能是否可以接受？我问这个看看你是否可能有一些同步问题。您执行多少计算？简单的总结只是一个例子？
您知道 SIMD 指令，对吧？它们可以同时对 2、4、8 或 16 个（或更多，取决于硬件）值执行相同的操作。这可能是时间相同的原因，因为它有效地完成了相同数量的工作，但使用的是 SIMD 指令。这在 GPU 中非常普遍，但在 CPU 上也可以观察到较小程度。 4 个浮点数对 SSE 寄存器有好处，听起来差不多。正如 Csaba 指出的那样，你也可能会成为内存传输的瓶颈。
嗯，我使用 float4 的原因是这样的：假设我们有 256 个线程可以同时运行。我们有一个 1024*1024 的图像。如果我们使用第一个内核，则 256 个线程首先处理前 256 列，然后处理接下来的 256 列，以此类推，直到所有列都处理完毕。在这种情况下，将消耗 4*t，其中 t 代表 256 个线程处理 256 列所消耗的时间。但是，如果我们使用第二个内核，由于每个线程处理图像的一行中的 4 个元素，所以256个线程可以处理1024列，时间为t。这只是一个估计。 @托马斯
我的意思是，在具有足够宽 SIMD 寄存器的硬件上，添加/乘以两个 float4 需要与添加/乘以两个浮点数相同的时间。因此，从这个意义上说，两个内核花费大约相同的时间这一事实应该不足为奇。
我不知道我的估计是否正确。实际上，我将 float4 用于我编写的另一个内核，它实际上实现了一些加速。所以我为这篇文章中提到的内核尝试了 float4。但我无法获得任何加速。

标签： opencl

【解决方案1】：

像 float4 这样的 OpenCL 矢量数据类型更适合较旧的 GPU 架构，尤其是 AMD 的 GPU。现代 GPU 没有可用于单个工作项的 SIMD 寄存器，它们在这方面是标量的。对于 NVIDIA Kepler GPU 和 Intel HD 集成显卡上的 OpenCL 驱动程序，CL_DEVICE_PREFERRED_VECTOR_WIDTH_* 等于 1。因此，在现代 GPU 上添加 float4 向量应该需要 4 次操作。另一方面，英特尔酷睿 CPU 上的 OpenCL 驱动程序的 CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT 等于 4，因此可以一步添加这些向量。

【讨论】：

【解决方案2】：

您正在直接从“src”数组（全局内存）中读取值。这通常比私有内存慢 400 倍。您的瓶颈肯定是内存访问，而不是“添加”操作本身。

当您从 float 移动到 float4 时，由于 GPU 能够处理矢量，矢量运算（加/乘/...）效率更高。但是，对全局内存的读/写保持不变。由于这是主要瓶颈，您根本看不到任何加速。

如果你想加快你的算法，你应该转移到本地内存。但是，您必须手动解决内存管理和正确的块大小。

【讨论】：

【解决方案3】：

您使用哪种架构？

使用 float4 具有更高的指令级并行度（然后需要少 4 倍的线程），因此理论上应该更快（参见 http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf）

但是我在你的内核中正确理解你正在做前缀和（你在每次迭代 y 时存储部分和）？如果是这样，由于存储，瓶颈在于内存写入。

【讨论】：

这是一个非常有趣的链接。我很惊讶 float4 有所作为。我没有想到 ILP。如果没有 ILP，我认为这无关紧要。线程数减少四倍，但每个线程的计算量却是四倍。

【解决方案4】：

我认为 GPU 上的 float4 不是 OpenCL 中的 SIMD 操作。换句话说，如果您添加两个 float4 值，则总和分四步完成，而不是一次完成。 Floatn 是真正为 CPU 设计的。在 GPU 上，浮动仅作为一种方便的语法，至少在 Nvidia 卡上。 GPU 上的每个线程都像没有 SIMD 的标量处理器一样工作。但是warp中的线程并不像它们在CPU上那样独立。考虑 GPGPU 模型的正确方法是单指令多线程 (SIMT)。 http://www.yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html

您是否尝试过在 CPU 上运行您的代码？我认为使用 float4 的代码应该比 CPU 上的标量代码运行得更快（可能快四倍）。此外，如果您有一个带 AVX 的 CPU，那么您应该尝试 float8。如果在 CPU 上 float4 代码比 float8 更快，那么在带有 AVX 的 CPU 上应该会更快。

【讨论】：

在遵循超标量模型的 nVidia GPU 下确实如此，但 AMD GPU 的 SIMD 寄存器非常宽（我上次检查过的为 8 或 16）。
这可能是真的，我不了解 AMD GPU，但它似乎与 GPGPU 型号不匹配。扭曲实际上是一个巨大的 SIMD 操作，所以它使用 SIMD，它只是被隐藏了。如果 AMD 每个线程都有 SIMD 寄存器，为什么他们不每个寄存器只有 8 或 16 个线程（这是 Nvidia 所做的）而不是一个线程和一个 8/16 SIMD 寄存器？
两者都有，iirc，warp（AMD 术语中的“波前”）架构有点不同，但它们仍然存在。 AMD GPU 没有“一个线程”：p 不过，我不知道他们的动机，但这就是为什么 AMD GPU 通常更适合处理令人尴尬的并行问题，它们只是对同等硬件具有更多的原始计算处理能力。 Nvidia GPU 还有其他优势。

【解决方案5】：

尝试为内核定义__属性__并查看运行时间的变化例如尝试定义：

__内核无效__属性__((vec_type_hint(int)))

或

__内核无效__属性__((vec_type_hint(int4)))

或者你想要的一些floatN

【讨论】：