糟糕的 OpenCL ImageSampling 性能与 OpenGL TextureSampling
Posted
技术标签:
【中文标题】糟糕的 OpenCL ImageSampling 性能与 OpenGL TextureSampling【英文标题】:Abysmal OpenCL ImageSampling performance vs OpenGL TextureSampling 【发布时间】:2012-06-14 13:58:48 【问题描述】:我最近将我的 volumeraycaster 从 OpenGL 移植到 OpenCL,这将 raycaster 的性能降低了大约 90%。我跟踪了 OpenCL 的图像采样函数的性能下降,这些函数比相应的 OpenGL 纹理采样函数慢得多。通过删除图像采样函数和纹理采样函数,两种光线投射器实现的速度大致相同。 为了在不同的硬件上轻松地对函数进行基准测试,并排除我其余 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 图像//纹理:这将 OpenCL 性能提高了 100%,尽管 OpenGL 性能根本没有改变。
-使用“最近”而不是“线性”插值:没有性能变化
这让我想知道: 我是否犯了一个非常愚蠢的错误,降低了 OpenCL 的性能? 为什么 OpenCL 采样性能如此之低,尽管它应该使用与 OpenGL 相同的纹理硬件? 为什么我的复杂三线性插值函数实现比它的硬件实现快? 如何提高 OpenCL 中的采样性能,以便获得与 OpenGL 中相同的速度?
【问题讨论】:
您是否尝试过使用 CL/GL 互操作? (khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/…) 我过去使用过它,因为我使用 OpenGL 进行了相当多的渲染以及使用 OpenCL 的计算。这可能不是您的最终解决方案 - 但它可能有助于阐明实际问题。 我感觉 OpenCL 并没有冒昧地在内存中调整你的纹理:c.f. en.wikipedia.org/wiki/Z-order_curve。这会导致非常糟糕的缓存访问性能,尤其是 3D 纹理。 我想在这里补充的另一点是,Nvidia 正在积极尝试压制 OpenCL 的使用以促进 CUDA(并因此锁定供应商)。我不会依赖他们的 OpenCL 实现优于在(可比的)英特尔或 AMD GPU 上运行的相同程序。 你为什么禁止编译器展开你的循环? forloop 将进行多次迭代,您的大部分时间将花在执行比较和跳转指令上。 IIRC 图形硬件没有分支预测,这会使其变慢? @user1449137 这个问题最初发布已经 2.5 年了。我很好奇:您或其他用户是否仍然看到 OpenGL 在较新的硬件上使用更新版本的 OpenCL 和 OpenGL 优于 OpenCL(对于此测试)? 【参考方案1】:我怀疑某些显卡上最新的 NVidia 驱动程序中的 OpenCL 存在一些问题。 Here 和 here 是关于这些的一些报告。尝试在其他家族的 GPU 上重复测试。
【讨论】:
以上是关于糟糕的 OpenCL ImageSampling 性能与 OpenGL TextureSampling的主要内容,如果未能解决你的问题,请参考以下文章