糟糕的OpenCL ImageSampling性能与OpenGL纹理采样

我最近将我的volumeraycaster从OpenGL移植到了OpenCL,这使得raycaster的性能降低了大约90%。 我跟踪了OpenCL的图像采样函数的性能下降,这些函数比相应的OpenGL纹理采样函数慢得多。 通过去除图像采样函数和纹理采样函数,两个raycaster实现具有相同的速度。 为了轻松实现不同硬件上的function,并排除其他RT代码中的一些愚蠢的错误,我写了一个小的基准,比较OpenCL采样速度和OpenGL采样速度,并在不同的机器上testing它, OpenCL仍然只有OpenGL性能的10%。

基准的OpenCL HostCode(至less是其中最重要的部分):

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,一些费米Tessla卡,所有这些在OpenCL和OpenGL中都有相同的糟糕performance

– 不同的纹理格式(字节而不是浮动),不同的访问模式和不同的纹理大小:不增加

– 使用缓冲区而不是数据的图像,在CL内核中使用自写的三线性内插函数进行采样:将OpenCL性能提高了约100%

– 使用2D图像//纹理而不是3D图像//纹理:尽pipeOpenGL的性能没有改变,但OpenCL的性能提高了100%。

– 使用“最接近”而不是“线性”插值:没有性能改变

这让我想知道:我做了一个非常愚蠢的错误,从而降低了OpenCL性能? 为什么OpenCL采样性能如此之低,尽pipe它应该使用与OpenGL相同的纹理硬件? 为什么我的复杂三线性插值函数实现比硬件实现更快? 我怎样才能提高OpenCL的采样性能,使我可以像OpenGL一样的速度?

我怀疑某些显卡上的最新NVidia驱动程序中存在OpenCL问题。 这里和这里是关于这些的一些报道。 尝试从另一个家庭重复testingGPU。