OpenCL图像采样性能与OpenGL纹理采样相比表现极差

55
我最近将我的体积光线投射器从OpenGL移植到OpenCL,结果性能下降了约90%。我发现性能下降是由于OpenCL的图像采样函数比相应的OpenGL纹理采样函数要慢得多。通过删除图像采样函数和纹理采样函数,两个光线投射器实现的速度大致相同。 为了方便在不同硬件上测试功能,并排除其余RT代码中的一些愚蠢错误,我编写了一个小型基准测试,比较了OpenCL采样速度与OpenGL采样速度,并在不同机器上进行了测试,但OpenCL仍然只有OpenGL性能的约10%。
基准测试的OpenCL HostCode(至少是最重要的部分):
void OGLWidget::OCLImageSampleTest()
{
    try
    {
    int size=8;
    float Values[4*size*size*size];
    cl::Kernel kernel=cl::Kernel(program,"ImageSampleTest",NULL);
    cl::ImageFormat FormatA(CL_RGBA,CL_FLOAT);
    cl::Image3D CLImage(CLcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR ,FormatA,size,size,size,0,0,Values,NULL);


    cl::ImageFormat FormatB(CL_RGBA,CL_UNSIGNED_INT8);
    cl::Image2D TempImage(CLcontext, CL_MEM_WRITE_ONLY,FormatB,1024,1024,0,NULL,NULL );


    kernel.setArg(0, CLImage);
    kernel.setArg(1, TempImage);



    cl::Sampler Samp;
    Samp() = clCreateSampler( CLcontext(), CL_TRUE, CL_ADDRESS_REPEAT, CL_FILTER_LINEAR, NULL);
    kernel.setArg(2, Samp);

    QTime BenchmarkTimer=QTime();
    BenchmarkTimer.start();

    cl::KernelFunctor func = kernel.bind(queue, cl::NDRange(1024,1024), cl::NDRange(32,32));
    func().wait();

    int Duration =  BenchmarkTimer.elapsed();
    printf("OCLImageSampleTest: %d ms \n", Duration);
    }
    catch (cl::Error& err)
      {
        std::cerr << "An OpenCL error occured, " << err.what()
                  << "\nError num of " << err.err() << "\n";
        return;
      }

}

OpenCL内核:

void kernel ImageSampleTest( read_only image3d_t CoordTexture, write_only image2d_t FrameBuffer, sampler_t smp)
{
int Screenx = get_global_id(0);
int Screeny = get_global_id(1);

int2 PositionOnScreen=(int2)(Screenx,Screeny) ;

float4 Testvec=(float4)(1,1,1,1);
for(int i=0; i< 2000; i++)
{
Testvec+= read_imagef(CoordTexture,smp, (float4)(0+0.00000001*i,0,0,0)); // i makes sure that the compiler doesn't unroll the loop
}

uint4 ToInt=(uint4)( (uint) (Testvec.x), (uint) (Testvec.y) ,(uint)(Testvec.z),1);
write_imageui (     FrameBuffer,  PositionOnScreen, ToInt ); 

}

OpenGL片段着色器用于全屏四边形,其具有与OpenCL内核工作项相同数量的片段:

#version 150
uniform sampler3D Tex;
out vec4 FragColor;

void main()
{
FragColor=vec4(0,0,0,0);
for(int i=0; i<2000; i++)
{
FragColor+= texture(Tex,vec3(0+0.00000001*i,0,0),0);
}
}

此外,我已经尝试了以下方法来提高性能:
-更改工作组大小:没有性能提升
-不同的硬件:280 GTX、580 GTX、一些Fermi Tessla卡,它们在OpenCL vs OpenGL方面的性能都非常差
-不同的纹理格式(字节而不是浮点数)、不同的访问模式和不同的纹理大小:没有增加
-使用缓冲区而不是图像进行数据处理,并在CL Kernel中使用自编的三线性插值函数进行采样:将OpenCL性能提高了约100%
-使用2D图像/纹理而不是3D图像/纹理:这将OpenCL性能提高了100%,尽管OpenGL性能没有变化。
-使用“nearest”而不是“linear”插值:没有性能变化
这让我想知道:
-我有没有犯一个非常愚蠢的错误,导致OpenCL性能降低?
-为什么OpenCL采样性能如此低,尽管它应该使用与OpenGL相同的纹理硬件?
-为什么我的复杂三线性插值函数实现比其硬件实现更快?
-如何提高OpenCL的采样性能,以便可以与OpenGL具有相同的速度?

2
你尝试过使用CL/GL互操作吗?(http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clCreateFromGLTexture3D.html)我过去曾经使用过它,因为我同时使用OpenGL进行渲染和OpenCL进行计算。这可能不是你的最终解决方案,但它可能有助于揭示实际问题。 - Ani
5
我感觉OpenCL没有在内存中自由地交换你的纹理(swizzling),请参考http://en.wikipedia.org/wiki/Z-order_curve。这将导致非常糟糕的缓存访问性能,尤其是对于3D纹理。 - v.oddou
2
我想在这里补充的另一点是,Nvidia正在积极试图扼杀OpenCL的使用,以推广CUDA(从而实现供应商锁定)。我不会指望他们的OpenCL实现能够胜过在(可比较的)Intel或AMD GPU上运行的同一程序。 - Ani
2
为什么你禁止编译器展开循环?for循环将会有许多迭代,并且你的大量时间将花费在执行比较和跳转指令上。如果我没记错,图形硬件没有分支预测,这将使其变慢? - Emily L.
2
@user1449137 这个问题最初发布已经过去了2.5年。我很好奇:在更新的OpenCL和OpenGL版本以及新硬件上,您或其他用户是否仍然看到OpenGL在这个测试中表现优于OpenCL? - NauticalMile
显示剩余5条评论
1个回答

3
我怀疑最新的NVidia驱动程序在某些显卡上的OpenCL存在问题。这里这里是一些相关报告。尝试在另一系列的GPU上重复测试。

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接