Abysmal OpenCL ImageSampling performance vs OpenGL TextureSampling

Question

Recientemente he portado mi volumeraycaster desde OpenGL a OpenCL, lo que disminuyó el rendimiento del raycaster en un 90 por ciento. Seguí la disminución del rendimiento a las funciones de muestreo de imagen de OpenCL, que son mucho más lentas que las funciones de muestreo de texturas OpenGL respectivas. Al eliminar las funciones de muestreo de imágenes y las funciones de muestreo de texturas, ambas implementaciones de raycaster tenían aproximadamente la misma velocidad. Para poder comparar fácilmente las funciones en diferentes hardware y excluir algunos errores tontos en el resto de mi código RTs, escribí un pequeño punto de referencia que compara la velocidad de muestreo OpenCL con la velocidad de muestreo OpenGL y lo probó en diferentes máquinas pero OpenCL todavía tenía solo un 10% del rendimiento de OpenGL.

del punto de referencia OpenCL HostCode (al menos la parte más importante de la misma):

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 Kernel:

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 para un quad pantalla completa que tiene la misma cantidad de fragmentos como el kernel OpenCL tiene elementos de trabajo:

#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); 
} 
} 

Además, ya he intentado th e siguiente para aumentar el rendimiento:

-cambiando tamaño de grupo de trabajo: no más de rendimiento

hardware -Diferentes: 280 GTX, 580 GTX, alguna tarjeta de Fermi tessla, todos ellos tenían el mismo rendimiento pésimo en OpenCL vs OpenGL

-Diferentes formatos de textura (bytes en lugar de los flotadores), diferentes patrones de acceso y diferentes tamaños de textura: no hay un aumento

-Utilizar un amortiguador en lugar de una imagen de los datos y una función de interpolación trilineal auto escrito para la toma de muestras en el Núcleo CL: Aumentó el rendimiento de OpenCL por aproximadamente 100%

-Uso de una imagen 2D // textura en lugar de una imagen 3D // textura: Esto aumentó el rendimiento de OpenCL en un 100% aunque el rendimiento de OpenGL no cambió en absoluto.

-Utilizar "más cercano" en lugar de la interpolación "lineal": No hay cambio en el rendimiento

Esto me dejó pensando: ¿Hice un error muy estúpido lo que disminuye el rendimiento OpenCL? ¿Por qué el rendimiento del muestreo OpenCL es muy bajo, aunque debería usar el mismo hardware de textura que OpenGL? ¿Por qué la implementación de mi compleja función de interpolación trilineal es más rápida que su implementación de hardware? ¿Cómo puedo aumentar el rendimiento del muestreo en OpenCL para que pueda tener la misma velocidad que en OpenGL?

Answer 1

Sospecho que hay algún problema con OpenCL en los últimos controladores de NVidia en algunas tarjetas de video. Here y here son algunos informes sobre eso. Intenta repetir la prueba en la GPU de otra familia.