我最近将我的体积光线投射器从OpenGL移植到OpenCL,结果性能下降了约90%。我发现性能下降是由于OpenCL的图像采样函数比相应的OpenGL纹理采样函数要慢得多。通过删除图像采样函数和纹理采样函数,两个光线投射器实现的速度大致相同。
为了方便在不同硬件上测试功能,并排除其余RT代码中的一些愚蠢错误,我编写了一个小型基准测试,比较了OpenCL采样速度与OpenGL采样速度,并在不同机器上进行了测试,但OpenCL仍然只有OpenGL性能的约10%。
基准测试的OpenCL HostCode(至少是最重要的部分):
此外,我已经尝试了以下方法来提高性能:
-更改工作组大小:没有性能提升
-不同的硬件:280 GTX、580 GTX、一些Fermi Tessla卡,它们在OpenCL vs OpenGL方面的性能都非常差
-不同的纹理格式(字节而不是浮点数)、不同的访问模式和不同的纹理大小:没有增加
-使用缓冲区而不是图像进行数据处理,并在CL Kernel中使用自编的三线性插值函数进行采样:将OpenCL性能提高了约100%
-使用2D图像/纹理而不是3D图像/纹理:这将OpenCL性能提高了100%,尽管OpenGL性能没有变化。
-使用“nearest”而不是“linear”插值:没有性能变化
这让我想知道:
-我有没有犯一个非常愚蠢的错误,导致OpenCL性能降低?
-为什么OpenCL采样性能如此低,尽管它应该使用与OpenGL相同的纹理硬件?
-为什么我的复杂三线性插值函数实现比其硬件实现更快?
-如何提高OpenCL的采样性能,以便可以与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片段着色器用于全屏四边形,其具有与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 vs OpenGL方面的性能都非常差
-不同的纹理格式(字节而不是浮点数)、不同的访问模式和不同的纹理大小:没有增加
-使用缓冲区而不是图像进行数据处理,并在CL Kernel中使用自编的三线性插值函数进行采样:将OpenCL性能提高了约100%
-使用2D图像/纹理而不是3D图像/纹理:这将OpenCL性能提高了100%,尽管OpenGL性能没有变化。
-使用“nearest”而不是“linear”插值:没有性能变化
这让我想知道:
-我有没有犯一个非常愚蠢的错误,导致OpenCL性能降低?
-为什么OpenCL采样性能如此低,尽管它应该使用与OpenGL相同的纹理硬件?
-为什么我的复杂三线性插值函数实现比其硬件实现更快?
-如何提高OpenCL的采样性能,以便可以与OpenGL具有相同的速度?