Geekly Articles each Day

CUDA：典型的なタスクのパフォーマンスの側面

計算アルゴリズムの実装のビデオカードへの転送を開始する前に、必要なパフォーマンスの向上を得るか、単に時間を失うかを検討する必要があります。 また、メーカーが数百GFLOPSを約束しているにも関わらず、現在の世代のカードには独自の問題があり、それは事前によく知られています。 理論を深く掘り下げることはせず、いくつかの重要な実用的なポイントを検討し、いくつかの有用な結論を定式化します。

CUDAがどのように機能するかを大まかに理解し 、 CUDA Toolkitの安定バージョンを既にダウンロードしていると仮定します。

Core Duo E8400のミドルエンドGTX460グラフィックスカードを苦しめます。

関数呼び出し

はい、何かを計算したい場合、カードで実行される関数を呼び出さずにはできません。 これを行うには、最も簡単なテスト関数を作成します。

__global__ void stubCUDA( unsigned short * output)
{
// the most valid function: yep, does nothing.
}

__global__指定子を使用すると、CPUから関数を呼び出すことでGPUで関数を実行できます。
cudaThreadSynchronize();
stubCUDA<<<GRID, THREADS>>>(0);
cudaThreadSynchronize();

すべての関数呼び出しはデフォルトで非同期であるため、呼び出された関数が完了するまで待機するにはcudaThreadSynchronize（）呼び出しが必要です。

このようなブロックをループで実行してみましょう。GRID= 160、THREADS = 96の場合、1秒あたり約15,000回の呼び出しを受け取ります。

まったく厚くないとだけ言っておきましょう。 何もしない最も単純な関数でさえ、0.7 msより速く実行することはできません。

最初の仮定は、ほとんどの時間がスレッドの同期に費やされ、非同期呼び出しははるかに高速に処理されることです（ただし、特定のタスクで使用する方がより具体的です）。

ご覧ください。 同期なしでは、関数を毎秒73100回実行することができました。 注目すべきは、結果はまったく印象的ではありません。

最後のテストでは、GRID = THREADS = 1で関数を実行してみましょう。これにより、カード内に多数のスレッドを作成するオーバーヘッドが解消されるはずです。 しかし、これはそうではありません。毎秒同じ73000-73500の呼び出しを受け取ります。

とても道徳的な：

外部メモリアクセス

有用なものを見つけるには、入力データと出力データが必要です。 これを行うには、ビデオカードとの間でデータが転送される速度を理解する必要があります。 次の関数を使用します。
cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice);

はい、CUDAは非同期データ送信ツールも提供しますが、将来のパフォーマンスは同期機能と変わりません。

大きなブロックをコピーします：cudaMemcpyHostToDeviceおよびcudaMemcpyDeviceToHostの指示のように、大きなブロック（100メガバイト以上）で約2 GB /秒のパフォーマンスが得られます。 一般的に、これは非常に良いです。

非常に小さな構造では事態はさらに悪化します。 4バイトを送信することにより、 1秒間に22,000コールを超えることはありません。 88 kb / s

道徳：

内部メモリアクセス

データをカードに転送した後、それらの操作を開始できます。 ビデオメモリへのアクセスのおおよその速度を評価したいと思います。 これを行うには、次の関数を作成します。
__global__ void accessTestCUDA(unsigned short * output, unsigned short * data, int blockcount, int blocksize)
{
// just for test of max access speed: does nothing useful
unsigned short temp;
for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
temp = data[j];
}
}
output[0] = temp;
}

GRIDおよびTHREADSパラメーターは、ここでその理由を説明するまでは既に使用されていますが、信じてください。すべてが正しいはずです。 うるさい人は、同期の欠如のために結果のスペルが間違っていると言いますが、私たちはそれを必要としません。

したがって、任意の読み取りで約42 GB /秒を取得します。 これはまったく悪くありません。

次に、入力データを出力にコピーするように関数を変更します。 意味はありませんが、ビデオメモリへの記録速度を評価することができます（変更は完全に単純なので、コードを複製しません）。

I / Oで約30 GB /秒を取得します。 悪くもありません。

実際には、シーケンシャルな（多少の違いはあるが）メモリアクセスを使用したことを修正する必要があります。 任意の数字の場合、最大2倍劣化する可能性がありますが、これは問題ではありませんか？

道徳：

算術演算

非常に単純な例を省略し、有用なことを行います。 つまり、画像の正規化（ピクセル[t]：=（ピクセル[t] -sub）*係数）。 実際のコード：
__global__ void normalizeCUDA(unsigned short * data, int blockcount, int blocksize, float sub, float factor)
{
for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
register float d = ( float )data[j];
d = (d - sub) * factor;
data[j] = (unsigned short )d;
}
}
}

ここでは、実数へのキャスト、ADDMUL、整数へのキャストという、一見コストのかかる3つの計算手順が使用されます。 フォーラムは、素材全体のキャストが非常に悪いことを恐れています。 おそらくこれは古い世代のカードには当てはまりましたが、今ではそうではありません。

合計処理速度： 26 GB /秒 。 3つの操作により、直接I / Oに比べてパフォーマンスがわずか13％低下しました。

コードを注意深く見ると、正規化は正しくありません。 整数に書き込む前に、実数を丸める必要があります。たとえば、round（）関数を使用します。 しかし、これをしないで、決して使用しないようにしてください！

ラウンド（d）： 20 GB / s 、さらにマイナス23％。
（符号なしショート）（d + 0.5）： 26 GB / s 、実際には測定誤差内の時間は変更されていません。

道徳：

論理演算

論理演算の速度を評価してみましょう。同時にもう1つ良いことを行います。配列内の最小値と最大値を見つけます。 通常、この段階は正規化の前に行われます（そして、このためだけに書かれました）が、逆になります。 彼はもっと難しいです。 作業コードは次のとおりです。
__global__ void getMinMaxCUDA(unsigned short * output, unsigned short * data, int blockcount, int blocksize)
{
__shared__ unsigned short sMins[MAX_THREADS];
__shared__ unsigned short sMaxs[MAX_THREADS];

sMins[threadIdx.x] = data[0];
sMaxs[threadIdx.x] = data[0];

for ( int i = blockIdx.x; i < blockcount; i += gridDim.x)
{
int vectorBase = i * blocksize;
int vectorEnd = vectorBase + blocksize;

for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
register unsigned short d = data[j];
if (d < sMins[threadIdx.x])
sMins[threadIdx.x] = d;
if (d > sMaxs[threadIdx.x])
sMaxs[threadIdx.x] = d;
}
}

__syncthreads();

if (threadIdx.x == 0)
{
register unsigned short min = sMins[0];
for ( int j = 1; j < blockDim.x; j++)
if (sMins[j] < min)
min = sMins[j];
if (min < output[0])
output[0] = min;
}

if (threadIdx.x == 1)
{
register unsigned short max = sMaxs[0];
for ( int j = 1; j < blockDim.x; j++)
if (sMaxs[j] > max)
max = sMaxs[j];
if (max > output[1])
output[1] = max;
}

__syncthreads();
}


ここでは、スレッドと共有メモリの同期なしでは実行できません。

合計速度： 29 GB / s 、正規化よりもさらに高速。

なぜ最小コードと最大コードを組み合わせたのですか-通常は両方が必要であり、呼び出しは個別に時間を失います（最初の段落を参照）。

一般に、条件付き操作ではビデオカードが悪いと言った人に石を投げます：人為的にこの断片をほぼ2倍遅くすることができましたが、このために条件の深さを4まで増やす必要がありました！ if（）if（）if（）if（）else if（）...

道徳：

複雑なデータ構造

アルゴリズムとデータ構造は強く結び付けられているという考えに基づいて（少なくともN. Wirthを思い出してください）、いくつかの複雑なデータ構造の状況を確認する必要があります。

ここで問題が発生します。データを関数に転送する場合、使用できるオブジェクトは2種類のみです-定数整数型（数字）とビデオメモリブロックへのリンク。

たとえば、リンクに基づいてツリーを構築するアイデアについては、すぐに説明します。

したがって、メモリの連続ブロックとこのブロックの要素への参照の配列という形で、複雑なデータ構造を表すことは変わりません。 したがって、ハッシュ配列、ツリー、および任意のデータ配列のインデックス構造を簡単に想像できます。

そのようなトリックの見返り-ダブルインデックスの必要性：
for ( int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
{
temp = data[index[j]+i];
}

このフラグメントは、インデックスとデータの内容とサイズに応じて、 10〜30 GB /秒の速度で動作します。 メモリ使用量の最適化を試みることができますが、最良の場合でもアクセス速度の25％を失います。 トリプルインデックスの動作はさらに悪く、パフォーマンスが40％〜60％低下します。

今日、私たちは多くを理解しました

ビデオカードの機能を適切に使用することにより、大量のデータがある場所、巧妙な計算の必要性、複雑なデータ構造の欠如など、画像処理、音声、ビデオなどのタスクで前例のないパフォーマンスを得ることができます。

トピックが気に入ったら、ビデオカード上のいくつかの便利なオブジェクトを計算する方法について説明します。距離マップ、画像の形態、検索インデックス、および十分に高速に動作し、同期で不要な問題を引き起こさない興味深いデータ構造を示します 。

Source: https://habr.com/ru/post/J119435/

More articles:


All Articles