9

GPGPUの「永続スレッド」についての論文をいくつか読んだことがありますが、よくわかりません。誰かが私に例を挙げたり、このプログラミング方法の使用法を教えてもらえますか？

「永続的なスレッド」を読んでグーグルした後、私が心に留めていること：

Presistent Threadsスレッドを実行し続け、多くの作業を計算するのは、whileループにすぎません。

これは正しいです？前もって感謝します

参照：http ：//www.idav.ucdavis.edu/publications/print_pub？pub_id = 1089 http://developer.download.nvidia.com/GTC/PDF/GTC2012/PresentationPDF/S0157-GTC2012-Persistent-Threads-Computing .pdf

4

2 に答える 2

11

CUDAは、単一命令複数データ（SIMD）プログラミングモデルを活用します。計算スレッドはブロックに編成され、スレッドブロックは別のストリーミングマルチプロセッサ（SM）に割り当てられます。SMでのスレッドブロックの実行は、スレッドをワープのスレッドに配置することによって実行され32ます。各ワープはロックステップで動作し、異なるデータに対してまったく同じ命令を実行します。

一般に、GPUをいっぱいにするために、カーネルは、実際にSMでホストできるはるかに多くのブロックで起動されます。すべてのブロックをSMでホストできるわけではないため、作業スケジューラは、ブロックの計算が終了したときにコンテキストスイッチを実行します。ブロックの切り替えはスケジューラーによって完全にハードウェアで管理され、プログラマーにはブロックがSMにスケジュールされる方法に影響を与える手段がないことに注意してください。これにより、SIMDプログラミングモデルに完全に適合せず、作業の不均衡が生じるすべてのアルゴリズムの制限が明らかになります。実際、ブロックの最後のスレッドの実行が終了するまで、ブロックは同じSM上のA別のブロックに置き換えられません。BA

CUDAはハードウェアスケジューラーをプログラマーに公開しませんが、永続スレッドスタイルはワークキューに依存することでハードウェアスケジューラーをバイパスします。ブロックが終了すると、キューでさらに作業がないかチェックし、作業がなくなるまでチェックを続けます。作業がなくなると、ブロックは終了します。このようにして、カーネルは、使用可能なSMの数と同じ数のブロックで起動されます。

永続スレッドの手法は、プレゼンテーションから抜粋した次の例でよりよく示されています。

「GPGPU」コンピューティングとCUDA/OpenCLプログラミングモデル

別のより詳細な例は、論文で利用可能です

GPUでのレイトラバーサルの効率を理解する

// Persistent thread: Run until work is done, processing multiple work per thread
// rather than just one. Terminates when no more work is available

// count represents the number of data to be processed

__global__  void persistent(int* ahead, int* bhead, int count, float* a, float* b)
{
    int local_input_data_index, local_output_data_index;
while ((local_input_data_index = read_and_increment(ahead)) <   count)
{                                   
        load_locally(a[local_input_data_index]);

        do_work_with_locally_loaded_data();

        int out_index = read_and_increment(bhead);

        write_result(b[out_index]);
    }
}

// Launch exactly enough threads to fill up machine (to achieve sufficient parallelism 
// and latency hiding)
persistent<<numBlocks,blockSize>>(ahead_addr, bhead_addr, total_count, A, B);
于 2014-06-10T17:06:53.367 に答える
4

非常に理解しやすいです。通常、各作業項目は少量の作業を処理しました。ワークグループの切り替え時間を節約したい場合は、ループを使用して1つのワークアイテムに多くの作業を処理させることができます。たとえば、1つの画像があり、それは1920x1080であり、1920の作業項目があり、各作業項目はループを使用して1080ピクセルの1つの列を処理します。

于 2014-01-20T07:45:45.897 に答える