1

VS 2012 および機能 3.5 (Titan および K20) で CUDA 5 を使用する。

カーネル実行の特定の段階で、生成されたデータ チャンクをホスト メモリに送信し、データの準備ができていることをホストに通知して、ホストがデータを処理できるようにしたいと考えています。

次の理由により、カーネルの実行が終了するまでデバイスからデータを読み取ることができません。

  1. データは、計算されるとデバイスに関係しなくなるため、最後まで保持しても意味がありません。
  2. データ サイズが大きすぎてデバイス メモリに収まらず、最後まで待機します。
  3. ホストは、データの処理を開始するために、カーネルの実行が終了するまで待つ必要はありません。

私がたどらなければならない道と、私の要件を達成するために使用しなければならない可能性のあるcudaの概念と機能を教えてもらえますか? 簡単に言えば、どうすればホストに書き込み、チャンク データがホスト処理の準備ができていることをホストに通知できますか?

注: 各スレッドは、生成されたデータを他のスレッドと共有せず、独立して実行されます。したがって、私の知る限り (間違っている場合は訂正してください)、ブロック、スレッド、およびワープの概念は質問に影響しません。言い換えれば、それらが答えに役立つ場合、私はそれらの組み合わせを自由に変更できます。

以下は、私がやろうとしていることを示すサンプルコードです:

#pragma once
#include <conio.h>
#include <cstdio>
#include <cuda_runtime_api.h>

__global__ void Kernel(size_t length, float* hResult) 
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // Processing multiple data chunks
    for(int i = 0;i < length;i++)
    {
        // Once this is assigned, I don't need it on the device anymore.
        hResult[i + (tid * length)] = i * 100;
    }

}

void main()
{
    size_t length = 10;
    size_t threads = 2;
    float* hResult;
    // An array that will hold all data from all threads
    cudaMallocHost((void**)&hResult, threads * length * sizeof(float));
    Kernel<<<threads,1>>>(length, hResult);
    // I DO NOT want to wait to the end and block to get the data
    cudaError_t error = cudaDeviceSynchronize();
    if (error != cudaSuccess) { throw error; }
    for(int i = 0;i < threads * length;i++)
    {
        printf("%f\n", hResult[i]);;
    }
    cudaFreeHost(hResult);
    system("pause");
}
4

1 に答える 1

2

デバイスの概要:

  • データをデバイスのグローバル メモリ (以前に で割り当てたもの) に書き込むか、ホスト メモリ (以前に で割り当てたもの) に直接書き込む必要がありますcudaMalloccudaHostAlloc
  • 次の手順の前にすべてのデータが書き込まれていることを確認するために、単一のスレッドブロックからこの領域へのすべてのデータ書き込みを実行することをお勧めします。
  • threadfence()次に、次の手順の前に (デバイスのグローバル メモリを使用している場合) またはthreadfence_system()呼び出し (ホスト メモリを使用している場合)を発行する必要があります。
  • 次に、デバイスのグローバル メモリまたはホスト メモリ内の特別な場所に書き込みます。これをメールボックスの場所と呼びましょう。データの準備ができていることを示す特定の値を使用します。
  • オプションで、別の threadfence または threadfence_system 呼び出しを発行します

ホスト上:

  • カーネルを起動する前に、ホストはメールボックスの場所をデフォルト値に設定する必要があります。
  • カーネルを起動した後、ホスト スレッドはメールボックスの場所を「ポーリング」して、データの準備ができていることを示す特定の値を探す必要があります。
  • データの準備ができていることを示す特定の値が表示されると、ホストはデータを消費できます。
  • オプションで、このプロセスを繰り返したい場合、主催者はメールボックスの場所をデフォルト値にリセットできます。デバイスは、データ ブロックを新しいデータで更新する前に、このデフォルト値を確認できます。

上記のプロセスを使用しても、データが複数のスレッドブロックから生成/作成されている場合は、暗黙のデバイス全体の同期が必要であることに注意してください。デバイス全体で利用できる唯一の単純な同期は、カーネルの起動 (具体的にはカーネルの完了) です。単一のスレッドブロックからデータをコピーすると、デバイス全体の同期の要件がこの特定のシーケンスから (このシーケンスの前のどこかに) 移動するだけです。

あなたが与えた理由は、コードをリファクタリングしてカーネル起動ごとにデータを作成することができなかったことを実際に示唆していません。これにより、これらの問題がきちんと解決され、上記のプロセスも不要になります。

編集:コメントの質問に応答します。カーネル呼び出しごとに 1 つのデータ チャンクを配信するようにコードをリファクタリングする方法について、具体的な例がなければ、より具体的に説明することは困難です。

グローバル メモリに 30 フレームのビデオ シーケンスが保存されている画像処理のケースを考えてみましょう。カーネルは、何らかのアルゴリズムに従って各フレームを処理し、処理されたデータをホストで利用できるようにします。

あなたの提案では、カーネルがフレームの処理を完了した後、データの準備ができていることをホストに通知し、次のフレームの処理に進むことができます。問題は、フレームが複数のスレッドブロックによって処理されている場合、すべてのスレッドブロックがそのフレームの処理をいつ完了したかを知る簡単な方法がないことです。デバイス全体の同期バリアが必要かもしれませんが、カーネル呼び出しメカニズムを介する場合を除いて、便利には存在しません。ただし、おそらくそのようなカーネル内では、次のようなシーケンスが存在する可能性があります。

  • while (more_frames)
    • フレームを処理する
    • シグナルホスト
    • フレーム ポインタをインクリメントする

リファクタリングされたアプローチでは、ループをカーネルの外に移動して、コードをホストします。

  • while (more_frames)
    • カーネルを呼び出してフレームを処理する
    • フレームを消費する
    • フレーム ポインタをインクリメントする

これを行うことで、カーネルは、フレーム処理がいつ完了したかを知るために必要な明示的な同期をマークし、データを消費できるようにします。

于 2013-05-06T10:16:21.003 に答える