Ужасная производительность OpenCL ImageSampling по сравнению с OpenGL TextureSampling

Недавно я перенес свой VolumeRaycaster с OpenGL на OpenCL, что снизило производительность Raycaster примерно на 90 процентов. Я отследил снижение производительности функций выборки изображений OpenCL, которые намного медленнее, чем соответствующие функции выборки текстур OpenGL. После удаления функций выборки изображений и функций выборки текстур обе реализации raycaster имели примерно одинаковую скорость. Чтобы легко протестировать функции на различном оборудовании и исключить некоторые глупые ошибки в остальной части моего кода RT, я написал небольшой тест, который сравнивает скорость выборки OpenCL со скоростью выборки OpenGL, и протестировал его на разных машинах, но OpenCL по-прежнему имел около 10 % производительности OpenGL.

Тестовый 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-изображения//текстуры: это увеличило производительность OpenCL на 100 %, хотя производительность OpenGL не изменилась. вообще.

-Использование «ближайшей» вместо «линейной» интерполяции: без изменения производительности

Это заставило меня задуматься: Я сделал очень глупую ошибку, которая снижает производительность OpenCL? Почему производительность выборки OpenCL такая низкая, хотя она должна использовать то же аппаратное обеспечение текстур, что и OpenGL? Почему моя сложная реализация функции трилинейной интерполяции работает быстрее, чем ее аппаратная реализация? Как я могу увеличить производительность выборки в OpenCL, чтобы иметь ту же скорость, что и в OpenGL?