与OpenGL TextureSampling相比，糟糕的OpenCL ImageSampling性能

我最近将我的volumeraycaster从OpenGL移植到OpenCL，这使得raycaster的性能降低了大约90％ . 我跟踪了OpenCL的图像采样功能的性能下降，这比相应的OpenGL纹理采样功能慢得多 . 通过删除图像采样功能和纹理采样功能，两个raycaster实现具有大致相同的速度 . 为了轻松地在不同的硬件上使用功能，并在其余的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 FragmentShader与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和OpenGL中具有相同的糟糕表现

• 不同的纹理格式（字节而不是浮点数），不同的访问模式和不同的纹理大小：没有增加

• 使用缓冲区而不是图像作为数据，并使用自编写的三线性插值函数进行CL内核中的采样：将OpenCL性能提高约100％

• 使用2D图像//纹理而不是3D图像//纹理：虽然OpenGL性能根本没有变化，但OpenCL性能提高了100％ .

• 使用“最近”而不是“线性”插值：没有性能变化

这让我想知道：我是否犯了一个非常愚蠢的错误，这会降低OpenCL的性能？为什么OpenCL采样性能如此之低，尽管它应该使用与OpenGL相同的纹理硬件？为什么我的复杂三线性插值函数实现比其硬件实现更快？如何在OpenCL中提高采样性能，以便与OpenGL具有相同的速度？