Abysmal OpenCL ImageSampling performance vs OpenGL TextureSampling
(1)
Recientemente transferí 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 benchmark 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.
El OpenCL HostCode del benchmark (al menos la parte más importante):
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;
}
}
Kernel 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 para un quad de pantalla completa que tiene la misma cantidad de fragmentos que 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 lo siguiente para aumentar el rendimiento:
-cambio de tamaño del grupo de trabajo: no aumenta el rendimiento
-Distintos Hardware: 280 GTX, 580 GTX, alguna tarjeta Fermi Tessla, todos tuvieron el mismo rendimiento abismal en OpenCL vs OpenGL
-Diferentes formatos de textura (bytes en lugar de flotantes), diferentes patrones de acceso y diferentes tamaños de texturas: sin aumento
-Utilizar un búfer en lugar de una imagen para los datos y una función de interpolación trilineal auto escrita para el muestreo en el Kernel CL: Aumentó el rendimiento de OpenCL en aproximadamente un 100%
-Uso de una textura // de imagen 2D en lugar de una textura // de imagen 3D: esto aumentó el rendimiento de OpenCL en un 100%, aunque el rendimiento de OpenGL no cambió en absoluto.
-Utilizar la interpolación "más cercana" en lugar de "lineal": sin cambio de rendimiento
Esto me dejó pensando: ¿cometí un error muy estúpido que disminuye el rendimiento de OpenCL? ¿Por qué el rendimiento de muestreo de OpenCL es muy bajo, aunque debería usar el mismo hardware de textura que OpenGL? ¿Por qué mi compleja implementación de la 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?