使用float4时加速,opencl

Question

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

Answer 1

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