4

我有以下 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。但是，无论图像有多大，两个内核几乎消耗相同的时间。我不知道为什么。你们能给我一些想法吗？吨

4

5 回答 5

7

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

于 2014-05-02T05:11:04.710 回答
4

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

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

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

于 2013-04-30T12:40:03.183 回答
2

您使用哪种架构？

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

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

于 2013-04-28T07:45:00.557 回答
0

我认为在 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。如果 float4 代码在 CPU 上比 float8 更快，那么在带有 AVX 的 CPU 上应该更快。

于 2013-04-28T07:43:35.700 回答
0

try to define __ attribute __ to kernel and see changes in run timing for example try to define:

__ kernel void __ attribute__((vec_type_hint(int)))

or

__ kernel void __ attribute__((vec_type_hint(int4)))

or some floatN as you want

read more: https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/functionQualifiers.html

于 2016-09-26T03:27:34.980 回答