2012-01-24 5 views
4

私は多くのBLAS gemv操作を行う関数を書いています。OpenCLまたはCUDAコールのオーバーヘッド?

私はGPUでこれを行うことができますし、私はcuBlasで試してみました。

私の問題は、私の行列とベクトルがかなり小さく、100x100行列と100ベクトルです。 CuBlasはCPUに比べて時間がかかり、なぜCPUの高速キャッシュとGPUへの呼び出しに大きなオーバーヘッドが混在するのかが分かります。

したがって、私はGPUに通話を通信するのにかかる時間を測定する賢明な方法を見つけようとしています。

これは、CUDAが呼び出しをセットアップしてグラフィックスプロセッサに送信するのにかかる時間です。これは、実際に行列 - ベクトル乗算を行う時間をカウントしません。

どうすればいいですか?

+0

CUDAにジョブを送信するのか、それとも興味のあるものだけを動的に選択できるようにするには? – Rup

+0

@Rup:コールが実際に費やしていることを理解することに興味があります。遅いコードが私のせいであるのか、それとも単に建築物の産物なのかを判断してください。 –

+1

少量のデータでは、あなたの頭上だけでなく、可能な並行性が欠けています。 GPUは待ち時間を隠すのに十分なスレッドを持っていることに熱心に依存しています(これはCPUのためにGPUではかなり悪いです)。呼び出しのオーバーヘッドがなくても、作業がスレッドの**ロット**に分割されていない限り、GPUはCPUより遅くなる可能性があります。たくさんのスレッドを簡単に意味することができます。 – Grizzly

答えて

8

更新:以下の結果は、(nVidiaの7800 GTX)2005ハードウェア上で手書きFFTのGPUのアルゴリズムのためのものであるが、CPU-GPUのtranferの原理を示しているが、オーバーヘッドがない

をボトルネックGPUプログラムをコンパイルし、GPUとホスト間でデータを転送します。 CPUはキャッシュ内で完全に実行できる機能に高度に最適化されており、DDR3メモリのレイテンシはGPUにサービスするPCI-Expressバスよりはるかに低くなっています。私はGPU FFTルーチン(CUDA以前)を書くときに自分自身を経験しました。 this related questionを参照してください。

 
N  FFTw (s) GPUFFT (s) GPUFFT MFLOPS GPUFFT Speedup 
8  0   0.00006  3.352705  0.006881 
16  0.000001 0.000065 7.882117  0.010217 
32  0.000001 0.000075 17.10887  0.014695 
64  0.000002 0.000085 36.080118  0.026744 
128  0.000004 0.000093 76.724324  0.040122 
256  0.000007 0.000107 153.739856  0.066754 
512  0.000015 0.000115 320.200892  0.134614 
1024 0.000034 0.000125 657.735381  0.270512 
2048 0.000076 0.000156 1155.151507  0.484331 
4096 0.000173 0.000215 1834.212989  0.804558 
8192 0.000483 0.00032  2664.042421  1.510011 
16384 0.001363 0.000605 3035.4551  2.255411 
32768 0.003168 0.00114  3450.455808  2.780041 
65536 0.008694 0.002464 3404.628083  3.528726 
131072 0.015363 0.005027 3545.850483  3.05604 
262144 0.033223 0.012513 3016.885246  2.655183 
524288 0.072918 0.025879 3079.443664  2.817667 
1048576 0.173043 0.076537 2192.056517  2.260904 
2097152 0.331553 0.157427 2238.01491  2.106081 
4194304 0.801544 0.430518 1715.573229  1.861814 

上記の表は、GPU FFT実装とCPU実装のタイミングをカーネルサイズに基づいて示しています。より小さいサイズでは、GPUとの間でのデータ転送が支配的です。より小さなカーネルはCPU上で実行でき、いくつかの実装/サイズは完全にキャッシュ内で実行できます。これにより、CPUは小規模な操作に最適です。

一方、GPUとの最小限の動きでデータの大量の作業を実行する必要がある場合、GPUはCPUを落とします。

あなたの例で効果を測定する限り、私は上記のような実験を行うことをお勧めします。行列のサイズごとに計算されたFLOPSを計算し、さまざまなサイズの行列に対してCPUとGPUでテストを実行してみましょう。 CSVファイルに、GPUとCPUのサイズ、時間、FLOPSを出力します。どのプロファイリングでも、数百回のコードを繰り返し実行し、すべての時間を計ってから、合計時間を反復で除算してループ時間を取得します。あなたのアルゴリズムが許せば(例えば、100x10ではなく10x100)、異なる形の行列を試してみてください。

このデータを使用すると、オーバーヘッドがどのようなものかを知ることができます。正確には、同じ実験を繰り返し実行しますが、GPUで実行される内部シェーダコードをノーオペレーション(単に入力から出力にコピー)に置き換えてください。このことができます

希望、

1

あなたは、イベントが、キューに入れられ提出され、開始され、あなたのバッファ転送イベントにclGetEventProfilingInfoを使用して仕上げたデバイスからナノ秒の時間を得ることができます。

詳細は、どのようにここでそれを設定するには:http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html

私は100×100の行列のために、あなたがクランチングのためにCPUに付着したほうが良いことだと思います。同時に多くのパケットを掛けなければ、(小さな)転送オーバーヘッドと通常ははるかに低いクロック速度のために、gpuの利点はほとんど目立たないでしょう。可能な限り多くのローカルデータを使用するようにカーネルを調整してください。ハードウェア上では、作業グループごとに32KBがあり、2つの100x100マトリックスを保持するのに十分なはずです。組み込みドットプロダクト関数も非常に便利です。彼らは、カーネルの最適化について詳細に話を http://developer.amd.com/afds/pages/OLD/sessions.aspx 、および最適なサイズをハードコード:

がADFS昨年のこのことについて素晴らしい講演がありました(2908セッションIDを参照してください)。

1

あなたのマトリックスはすでにGPUに入っていますか? そうでない場合、CUBLASはそれらをあなたのために転送することがあります(サンクと呼ばれます)。これは追加のオーバーヘッドです。

また、GPUはそのような小さな計算では輝きません。つまり、結果を戻す必要があるため、おそらくCPUよりも遅くなります。 可能であれば、より大きな行列を使用してください。 それ以外の場合は、ストリーム(cudaStream_t)を使用して、GPUで複数の並列計算を開始することができます。このようなイベントで

あなたはCUDAカーネルの実行時間を計測したい場合、あなたはその(またはGPU上で計算何かを)囲む必要があり、CUDAランタイムAPIを使用して:

cudaEvent_t start, stop; 

cudaEventRecord(&start); 

struct timeval cpuStart, cpuEnd; 

gettimeofday(&cpuStart, 0); // get start time on CPU 

// Do something with CUDA on the GPU, e.g. call kernels, transfer memory, ... 

gettimeofday(&cpuEnd, 0); // get end time on CPU 

double seconds = cpuEnd.tv_sec - cpuStart.tv_sec; 
double microseconds = cpuEnd.tv_usec - cpuStart.tv_usec; 
double cpuDuration = (seconds * 1.0e6 + microseconds)/1.0e3; // in milliseconds 

cudaEventRecord(&stop); 

// Wait until the stop event occurred 
cudaError_t eventResult; 

do 
{ 
    eventResult = cudaEventQuery(stop); 
} 
while (eventResult == cudaErrorNotReady); 

// Assert there was no error; check the CUDA Toolkit Reference for further info 
assert(cudaSuccess == eventResult); // requires #include <assert.h> or <cassert> 

// Retrieve the time 
float gpuDuration = 0.0; // in milliseconds 
cudaEventElapsedTime(&gpuDuration, start, stop); 

// Release the event objects 
cudaEventDestroy(stop); 
cudaEventDestroy(start); 

前回の呼び出しでエラーが発生したため、CUDAのすべての呼び出しのエラーコードを確認することができます(前回の呼び出しからエラーが発生する可能性があるため)...

CUDAドライバAPIは、そのままでは動作しない可能性があります。申し訳ありません)。

EDIT:ちょうどカーネルの持続時間ではなく、呼び出し自体を測定したいと思っていました。 これは、コールのCPU時間を測定することで可能です。上記の更新されたコードを参照してください。 gettimeofdayはWindows(AFAIK)では使用できないため、これはLinuxのみで動作します。

+1

Windowsでは[QueryPerformanceCounter](http:// msdn。マイクロソフトのWebサイト(http://www.microsoft.com/japan/technet/itsolutions/ms644904)または[GetSystemTime](http://msdn.microsoft.com/en-us/library/windows/desktop/ms725473.aspx) – Rup

+0

デバイス上のデータであり、単純に Ax-> yを実行し、デバイス上にyを保持する必要があります。 –

+1

その場合、cublasDgemm()呼び出しの周りにgettimeofday()(またはWindows上の同様のメソッド)を置くことによって、CUBLASが実際のカーネルを起動する必要がある時間を測定できます。 私は自分で試していませんが、Parallel Nsight(Windowsの場合)またはVisual Compute Profiler(Linuxの場合はツールキットに含まれています)を使用して調べることができます。 私はそれを見つけることができませんが、私はCUDA 4のプロファイリングフックに関する何かを見てきたと確信しています... EDIT:CUDAのプロファイリングに関する興味深い情報があるこのPDFが見つかりました:http:// bit .ly/zn6jbP –

1

コールオーバーヘッドを見つけるには、できるだけ少ないCUDAカーネルを呼び出します。

for (int i=0; i<NLoops; i++) { 
    gettimeofday(&cpuStart, 0); // get start time on CPU 

    // Call minimal CUDA kernel 

    gettimeofday(&cpuEnd, 0); // get end time on CPU 

    // save elapsed time 
} 

上記のAlex P.のコードに従ってください。

カーネルで行う処理が少ないほど、呼び出しのオーバーヘッドだけが時間差になります。

NLoops(おそらく1,000,000)に対して適切な値を見つけるために少し実験を行います。経過時間がタイマーの間隔よりも長いことを確認してください。そうしないと、すべてがゼロになります。それが起こった場合、予測できる一定の時間間隔で実行するカーネルコードを書きます(それぞれxサイクルのn回のループ)。

cpuStartとcpuEnd(割り込み処理のような)の間で発生する可能性のある非CUDA計算をすべて削除するのは難しいですが、いくつかの実行と平均化を行うと良い結果が得られます。

関連する問題