Abysmal OpenCL ImageSampling 성능 대 OpenGL TextureSampling

Abysmal OpenCL ImageSampling 성능 대 OpenGL TextureSampling

---

저는 최근에 제 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 ); 

}

OpenCL 커널에 작업 항목이있는 것과 동일한 양의 조각이있는 전체 화면 쿼드 용 OpenGL FragmentShader :

#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 % 향상

-Using a 2D image//texture instead of a 3D image//texture:This increased the OpenCL performance by 100 % although the OpenGL performance didnt change at all.

-Using "nearest" instead of "linear" interpolation: No performance change

This left me wondering: Did I do a very stupid mistake which decreases the OpenCL performance? Why is the OpenCL sampling performance so low, although it should use the same texture hardware as OpenGL? Why is my complex trilinear interpolation function implementation faster than its hardware implementation? How can I increase the sampling performance in OpenCL so that i can have the same speed as in OpenGL?

---

I suspect there is some issue with OpenCL in latest NVidia drivers on some video cards. Here and here are some reports about those. Try to repeat test on GPU from another family.

ReferenceURL : https://stackoverflow.com/questions/10983906/abysmal-opencl-imagesampling-performance-vs-opengl-texturesampling