55

私は最近、自分の volumeraycaster を OpenGL から OpenCL に移植しました。これにより、raycaster のパフォーマンスが約 90% 低下しました。それぞれの OpenGL テクスチャサンプリング関数よりもはるかに遅い、OpenCL のイメージサンプリング関数のパフォーマンス低下を追跡しました。画像サンプリング関数とテクスチャ サンプリング関数を削除することで、両方の raycaster 実装の速度はほぼ同じになりました。さまざまなハードウェアで関数を簡単にベンチングし、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 で同じ最悪のパフォーマンスを示した

- さまざまなテクスチャ フォーマット (float ではなくバイト)、さまざまなアクセス パターン、さまざまなテクスチャ サイズ: 増加なし

- データ用の画像の代わりにバッファを使用し、CL カーネルでのサンプリング用に自作のトライリニア補間関数を使用: OpenCL のパフォーマンスを約 100% 向上

-3D 画像//テクスチャの代わりに 2D 画像//テクスチャを使用: これにより、OpenCL のパフォーマンスは 100% 向上しましたが、OpenGL のパフォーマンスはまったく変化しませんでした。

- "linear" 補間の代わりに "nearest" 補間を使用: パフォーマンスの変更なし

これは私に疑問を投げかけました: OpenCL のパフォーマンスを低下させる非常にばかげた間違いをしましたか? OpenGL と同じテクスチャ ハードウェアを使用する必要があるのに、OpenCL のサンプリング パフォーマンスが非常に低いのはなぜですか? 複雑なトライリニア補間関数の実装がハードウェアの実装よりも速いのはなぜですか? OpenCL でのサンプリング パフォーマンスを向上させて、OpenGL と同じ速度にするにはどうすればよいですか?

4

1 に答える 1

3

一部のビデオ カードの最新の NVidia ドライバーでは、OpenCL に問題があると思われます。 こことここに、それらに関するいくつかのレポートがあります。別のファミリの GPU でテストを繰り返してみてください。

于 2013-11-04T07:55:01.480 に答える