記事では、 前の記事で述べたよりもはるかに高速に、バイナリラスタ上の接続オブジェクトを非常に迅速にリストする方法を説明します。 そのような速度がどこにあるように見えるでしょう。 ここで、数十ミリ秒の間、4096 x 4096の写真で「分解」します。 そして、問題自体は興味深いものですが、その解決策はかなり幅広い適用性を持つかなりシンプルで独創的な方法に基づいており、その主なテーマは「より簡単にし、何が起こるかを見ていきます」です。 この場合、CUDAをメインの計算機として使用しますが、「非常にシンプル」にしたいため、詳細はありません。
挑戦する
バイナリラスタ上のすべての接続オブジェクトを選択し（ピクセルの色は白黒のみ）、何らかの方法でそれらをマークする必要があります。 明確にするために、異なる色でマークします。 もちろん、「色-オブジェクトインデックス」の対応は非常に簡単です。 この例では、2つのポイントが垂直または水平に直接隣接している場合は接続されていますが（4接続）、アルゴリズムは8接続の場合に簡単に変更できます（これは主なことではありません）。 タスクはこの記事で非常に明確に述べられています 。
たとえば、元の画像：
このアルゴリズムは非常に大きな画像（最大16384x16384）で機能しますが、明白な理由により、記事のテキストでは縮小画像が使用されます。
探索するかどうか？
新しいタスクを受け取ったときに、どの解決方法が最適であるか、解決方法にどのような制限が生じるか、ソリューションの面倒さを考えることは非常に合理的です。 方法を選択する際に客観的になるには、多くのオプション、時には「多すぎる」を考慮する必要があります。そのため、主なことは、不要なオプションをすばやく排除するために正しく議論を開始することです。
最悪の「ベストケース」
明らかに、問題を解決するには、画像の各ピクセルを考慮する必要があります。 モデルコードがサイズ8192x8192の画像で実行される時間を見てみましょう。
for ( int x = 0; x < dims[0]; x++)
{
for ( int y = 0; y < dims[1]; y++)
{
image[x + dims[0]*y] = 0;
}
}
はい、サイクルを1つに折り畳むことで大幅に「最適化」できますが、合理的な方法では、画像の個々のピクセルにアクセスできる可能性があるため、インデックスは意図的に残されます。 「平均的な」プロセッサー（私の場合、Core 2 Duo E8400）で実行し、約680-780 msのランタイムを取得します。
大量か少量かを判断することは困難ですが、問題をより迅速に解決することはできないようです。
もしくはそうですね。 「すべての手段が優れている」ことと、問題の解決を大幅に簡素化し、より良い結果をもたらすテクノロジーがあるためです。 この場合、メモリを集中的に使用するタスクに対して非常に優れたパフォーマンスを発揮するCUDAテクノロジーが非常に役立ちます。
私の平凡なGTX460グラフィックスカードでは、約40 GB /秒のビデオメモリにアクセスできますが、メインメモリでは2 GB /秒です。 この処理速度は、大規模並列処理とかなり高速なメモリを使用して実現されます。
分割して征服する
ビデオカードでの計算の詳細を考慮すると、タスクの分離は非常に重要であることがわかります。スタックがなく、ストリーム間の同期が非常に速くないため、複雑なデータ構造を構築できません。
元の画像を分割する合理的な方法は、それを小さな「正方形」に分割することであるように思われます。それぞれが別々のストリームで処理されます。 これはCUDA計算モデルにうまく適合します。 また、過度の特異性を省略すると（1つの変換によって行われます）：
UID = threadIdx.x + blockDim.x * blockIdx.x
次に、元のタスクが、フローのすべての一意の識別子（UID）のサブタスクの結果の結合として表されることを取得します。 はい、私たちはまだ何をするのか正確にはわかりませんが、タスクを分解する方法はすでに知っています。
（色の暴動は作者のスタイルであり、グラデーションはすべてのUIDが異なることを示しています）元の画像をカバーできます：
各ストリームは、対応するブロック内のピクセル（たとえば、サイズが32 x 32）を処理します。その後、結果の同期と組み合わせが必要になる場合があります。
それと同じくらい簡単
数学では、1つの素晴らしい原則がよく使用されます。「何かが不可能であるが、本当にしたいなら、それができるということです」、実際には「後で」複雑さを「押し戻す」だけですが、この場合は便利になります。
最も単純な塗りつぶしアルゴリズムを採用し、各ブロック内で、境界効果をまったく考えずに、対応するポイント（まだ塗りつぶされていないそれぞれから開始）に適用します。
実際、ブロック内の画像ポイントを一意の識別子で埋めます。結果は予測可能です。
これまで、同じ「色」を持つ同じUIDの新しい塗りつぶしの繰り返しを開始しました。 しかし、これはほとんど意味がありません-このアプローチで（元の写真の）接続された領域が1つの色で塗りつぶされることを保証できない場合、少なくとも切断された領域を異なる「色」で塗りつぶすことができます。 したがって、UIDから取得した新しい色と、全体の画像でこの色の一意性を保証する特別な添加剤を使用して、塗りの新しい反復を実行します（スレッドの総数に反復数を掛けた値を追加するだけです）。
写真で処理された正方形の境界を表す「グリッド」で勝ったことを理解するのは簡単です。
これで、各ボックス内で、接続された領域は1色で塗りつぶされ、切断された領域は異なります。 これはO（n ^ 3）（はい、塗りつぶしアルゴリズムの最悪の実装を選択しました）ごとに機能します。各ボックスには（imageSize.x / n）*（imageSize.y / n）があり、そのほとんどは同時に十分に処理されますスレッド数（通常は約256）。
延期された困難
結果を1つにまとめることは残っています。 認識レベルでは、サブタスクに従来の方法を適用するよりも難しいという事実にもかかわらず、アルゴリズム的にははるかに高速です。
実際、どの接続色が同じ接続領域に属しているかを理解する必要があります。
これを行うには、グリッドのピクセル（図では正確に白で描かれているピクセル）を調べて、異なる色のピクセルがすぐ隣にあるかどうかを確認します。 そうである場合、両方の色は本質的に1つのオブジェクトを意味し、簡単に「塗り直し」、またはこの事実を「念頭に置いて」保持できます（通常はオブジェクトのパラメーターを取得するのに十分です）。
単純な算術は、グリッド内のピクセル数が画像の周囲に線形に依存するという事実につながります-正方形内のすべてのポイントを処理する必要はありませんが、それらの境界のみが必要です。
これらのポイントを通過した後、カラーマッチングマップ（c 1 == c 2 == c 3 ...）を作成し、結果を正確に描画できます。
実装と速度の見積もり
驚くべきことに、この方法の「ボトルネック」は、メインメモリからビデオメモリへのデータの転送です。ここには代替手段がありません。
START_ACTION( "Allocating GPU memory: image... " );
unsigned int * data_cuda;
cutilSafeCall( cudaMalloc(( void **)&data_cuda, data_cuda_size) );
END_ACTION(data_cuda_size);
START_ACTION( "Copying input data to GPU memory: image... " );
cutilSafeCall( cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice) );
END_ACTION(data_cuda_size);
結果の2 GB / sは驚くことではありません-これはRAMへのアクセス速度です。 使用する画像の場合、これらはミリ秒単位であり、怖いものではありません。
カーネルの呼び出し（ビデオカードで実行される機能）：
START_ACTION( "Calling CUDA kernel... " );
cutilSafeCall( cudaThreadSynchronize() );
processKernelCUDA<<<GRID, THREADS>>>(data_cuda, dims[0], dims[1]);
cutilSafeCall( cudaThreadSynchronize() );
END_ACTION(data_cuda_size);
結論はそれ自身を物語っています：
Calling CUDA kernel... 26.377720 msecs.
Throughput = 2.36942 GBytes/s
実際、ビデオメモリへの必要なデータ（画像）の転送が少ない時間に、すべてのサブタスクの実行を実現することができました。 これは、最悪の塗りつぶしアルゴリズムを選択し、その実装を絶対的に最適ではないという事実にもかかわらずです。 何かを最適化するのは理にかなっていますか？ あなたが決める:)
カーネルコードがなければ、ステートメントはおそらく不完全なままなので、愛してください：
__global__ void processKernelCUDA(unsigned int * image, int dim_x, int dim_y)
{
// Full task size: image dimensions
int TASK_SIZE = dim_x * dim_y;
// Max UID for full task (really count of 'quads')
int MAX_UID = TASK_SIZE / QUAD_SIZE / QUAD_SIZE;
// calculate UID by CUDA implicit parameters; WORKERS = GRID*THREADS
for ( int UID = threadIdx.x + blockDim.x * blockIdx.x; UID < MAX_UID; UID += WORKERS)
{
// count of quads for x-dimension
int quads_x = dim_x / QUAD_SIZE;
// current quad index form UID
int curr_quad_x = UID % quads_x;
int curr_quad_y = UID / quads_x;
// offset in global task
int offset = curr_quad_x * QUAD_SIZE + dim_x * curr_quad_y * QUAD_SIZE;
// id of 'color' to mark block, have to be >= 2
int COLOR_ID = 2 + UID;
do
{
// means that fill is not started yet
bool first_shot = true ;
// process quad starting with specified 'offset'
bool done_something;
do // fill step
{
done_something = false ;
for ( int local_y = 0; local_y < QUAD_SIZE; local_y++) // step by y
{
// calcluate point index in whole image
int start_x = offset + local_y * dim_x;
for ( int pos = start_x; pos < start_x + QUAD_SIZE; pos++) // step by x
{
if (image[pos] == 1) // is not marked with ID, is not blank
{
if ( first_shot /* start fill */ // else check neighbors:
|| image[pos-1] == COLOR_ID // x--
|| image[pos+1] == COLOR_ID // x++
|| local_y > 0 && image[pos-dim_x] == COLOR_ID // y--
|| local_y < QUAD_SIZE && image[pos+dim_x] == COLOR_ID ) // y++
{
image[pos] = COLOR_ID; // mark with ID
first_shot = false ; // fill already started
done_something = true ; // fill step does something
}
}
}
}
} while (done_something); // break if fill have no results
if (first_shot)
{
// no more work in this block, exit
break ;
}
else
{
// next iteration: restart block with different color ID
COLOR_ID += MAX_UID;
}
} while ( true );
}
}
入力は、画像へのポインタとXおよびYの画像の解像度です。内部では、サブタスク識別子が計算され、サブタスクに対応する正方形のオフセットが選択され、非再帰的な塗りつぶしアルゴリズムで塗りつぶされます。
成功の秘
もちろん、データを受信する時間よりも短い時間で解決されるタスクはそれほど多くありませんが、この場合は幸運でした。 そして運以外に-解決するテクノロジーの正しい選択と推論の正しい方法。これにより、かなり深刻なタスクの結果（種類）を1日以内に取得できます。
PS大きなリクエストです。写真をコピーしたり、許可なく再公開したりしないでください（このメソッドの著者です）。