3

これは、並列化しようとしている疑似コード (word2vec C コードから取得) です。まず、データ構造とそれに対応するサイズをリストし、次に疑似コードをリストします。

1.  long long sen[MAX_SENTENCE_LENGTH]  
// In the C code, MAX_SENTENCE_LENGTH = 1000. Increasing this should be  
//fine.

2.  float neu1[N] (hidden layer values)
//N is the length of each vector. For now, max N = 400

3.  float neu1e[N] (hidden layer error values)

4.  float syn0[V * N] (input to hidden layer weight matrix)
// For now, we can assume that V * N is small enough to be stored on the GPU
   // In the test data, V = 72k words

5.  float syn1neg[V * N] (back propagation weights used during negative  
sampling)

6. float exptable[1000] 

プログラムへの入力はテキスト ファイルです。次に、プログラムはそれを一度に 1 語ずつ処理して語彙を構築します。たとえば、テキスト ファイルに次の文があるとします。

「並列プログラミングは非常に興味深い」</p>

すると、語彙は次のようになります (コードは単語の頻度に基づいて語彙を並べ替えるため)。

            {“Very:2”, “Parallel:1”, “programming:1”, “is:1”,    “interesting:1”}
                   0      1               2              3                4

語彙を構築した後、コードは一度に 1000 語ずつテキストの処理を再開します。最初の 1000 語が に保存されsen[MAX_SENTENCE_LENGTH]、次にニューラル ネットワークが 内のすべての語についてトレーニングされsen、このプロセスがファイルの最後に到達するまで続きます。上記の文の場合、次のようにsenなります[1,2,3,0,0,4]

トレーニングが 1 回の反復でのみ行われると仮定すると、擬似コードは次のようになります。

for sen in text
{ 
    for word in sen
    {

        for (c = 0; c < N; c++) 
            neu1[c] = 0;

        for (c = 0; c < N; c++) 
            neu1e[c] = 0;   

       /*The variable window is a user supplied parameter. 
        It is used to consider the context  around a word in a sentence. 
        For example, if I am looking at the first word in the sentence
        (target word is word1), and window = 5, then the words in the 
        window = {word2, word3, word4, word5}. 
        If I am looking at the third word in the sentence 
        (target word is word3), then window = {word1, word2, word4, word5}*/    

        for word in window
        {
            for (c = 0; c < N; c++) 
            neu1[c] += syn0[c + word * N];
        }

        for (c = 0; c < N; c++) 
            neu1[c] /= window;

        //negative: number of negative samples to provide (assume it to be 
             //between 5 to 25)
        for (d = 0; d < negative + 1; d++) 
        {

            target = sen[random_index]  
            l2 = target * N;
            f = 0;
            for (c = 0; c < N; c++) 
            f += neu1[c] * syn1neg[c + l2];

           gradient = exptable[function of f] //f is calculated in the loop above

           for (c = 0; c < N; c++) 
              neu1e[c] += gradient * syn1neg[c + l2];

           for (c = 0; c < N; c++) 
              syn1neg[c + l2] += gradient * neu1[c];

          } //Negative Sampling ends    

        for word in window
        {
             for (c = 0; c < N; c++) 
                syn0[c + word * N] += neu1e[c];
        }

   } // word in sen loop ends

 } // sen in text loop ends

これを並列化する最善の方法は、文中の単語を並列処理することだと考えています。Nすべてのループを考慮すると、単一のスレッドがループごとに 1 回だけグローバル メモリ ( ) にアクセスするように、ワードごとにスレッドを使用する必要があると思いsyn0, syn1negます。また、neu1とのneu1e更新はすべて独立しているため、スレッドのプライベート メモリに常駐し、個別に更新することができます。

現在の主な懸念事項は次のとおりです。

  1. グローバル メモリ アクセスは、変数の値 (ボキャブラリのインデックス) に基づいてアクセスされるためsyn0、ランダムな方法で発生しています。そして、ご覧のとおり、文中の単語は順不同です。syn1negword

これは大きな問題ですか?または、GPU に十分な数のスレッドを与えることで、メモリのレイテンシを隠すことはできますか? また、N スレッド/ワードが syn0 および syn1neg のシーケンシャル データにアクセスするため、このアクセス パターンが実際にランダムかどうかはわかりませんが、N スレッドの次のセットはメモリ内の遠くにあるシーケンシャル データにアクセスする可能性があります。

  1. ネガティブ サンプリング ループでは、リダクション操作を実行する必要があります。変数fは内積の合計です。問題は、グローバルメモリにあるのneu1に対し、各スレッドのプライベートメモリに格納することを計画していることです。syn1neg

ネガティブ サンプリングには別のカーネルが必要ですか? Nスレッド/ワードを起動するだけではなく、別のアプローチが必要なようですが、どのアプローチが最適かはわかりません。

これらの懸念とは別に、このコードへのアプローチ方法に問題があるかどうかを提案してください。

4

1 に答える 1

-2

プロローグ:あなたはワームの缶を開けました (たとえSLOCプレゼントがなくても) では、複雑な主題から「逃げる」のではなく、象のスライスが全体として複雑な主題に対処する唯一のアプローチであるように思われるため、パートごとに提示された発言を受け入れることができることを願っています。プリンシパルの問題は、実装ドメインの個々のコード行のコンフォートゾーンに下ります。そこでは、アプリオリにまだ失われていない場合でも、全体像が通常見落とされます。


A1:はい、これは主要な事実です (別名「問題」 )。

GPU- デバイスは単一の命令としてシリコンで設計および最適化されていますSIMD-複数のデータハードウェアアーキテクチャであるため、コード+データ レイアウトの両方で最高のパフォーマンスを発揮します。 ) 実際、コアのオンチップ メモリ領域に収まる小さなメモリ領域 (キロバイト) ( -s スピルオーバーなし、例: LRU が維持する L1-キャッシュ ) したがって、アクセスごとの「理想化された」パフォーマンス壊滅的なレイテンシ ペナルティを導入しません。SIMD SMGPU-SM-REGISTERabout 350-700 nsgloMEM

[B]-Fig.2状態:あたり8 [ ] コア、あたり[ ]

TESLAのペア、マルチスレッド フェッチ & イシュー [ ] TPC (テクスチャ/プロセッサ クラスター) 1 つにつき16KBバンク、それぞれにつき 1 つの読み取り専用(高速/低合体衝突) TECH.GPU: NVIDIA CUDA C プログラミング ガイド、[PG-02829-001_v7.0]; 2015/03SM
SMXSM
SFUSM
MTIper
shaMEMSM
conL1cache SM
[B]

これは、ラスター画像処理 (小さなタイルで編成された行列畳み込み計算スタイルでデータの 2D 配列を処理する) でうまく機能し、同時に (確かにwarpSize-segmented )すべてのスレッドが同じSIMD-instructionを実行します (理想的には) 衝突しないデータ セル -- これがGPGPU最適です。

これはまた、そのような一意に段階的に漸進的で完全に調整された -ops を許可しない実際の操作は、SIMDパフォーマンスをブロックすることを意味します (スレッドは、(再) 同期のために、スレッド間の発散を待つことができます)。データ配信が最終的に行われるまでのリモートメモリアクセスのための障壁 & したがって、レイテンシーマスキングはますます実際にこれらの自然な障害を隠してPARALLEL、真のコード実行の錯覚を見ないようにします)。


A2:いいえ、ほとんどありません。その影響を定量的に評価し、結果として生じる実行時間の範囲の限界を実証するために利用できる、その場でのベンチマーク手法と証明。

kernel-deployment ディレクティブからいくつかのヘルプがありますが__launch_bounds__():

__global__ void
__launch_bounds__( 1,     // dim3tbGridSIZE <maxThreadsPerBLOCK>         COMPILE-TIME_ADVICE_FOR_OPTIMISING_COMPILER_____________REGISTERs, CACHE_, FETCH_, PROXIMITY_PATTERNs ANALYSES
                   1  /*, // dim3tBlockSIZE <minBlocksPerMULTIPROCESSOR> COMPILE-TIME_ADVICE_FOR_OPTIMISING_COMPILER_____________OPTIMUM_SCHEDULE_TO_FILL_FETCH_LATENCIES
                   ?,     // iAsyncSeqOfCmdsQUEUE_Stream_ID <<- TO LET BE FREELY ASSIGNABLE ... NON-BLOCKING EXEC'd KERNEL
                   0  */  // iSharedMemSIZE
                   )
                 Device_printf_GPU_CLK( int const iTag ){
                        ...
                        return;
}

広範な(ブルートフォーススキャンなど)「最適化」、むしろ機械的な調整、さまざまな起動パラメータ化のためのカーネルコンパイル(threading-3D-「geometry」)、コードの一般的な仮定への影響について公開された多くの論文があります-結果は常にカーネル固有であるため、設計を過大評価してはなりません (そして、実際には、アクセス階層SMX内のシリコンに限られたリソースの展開中にどの 3D ジオメトリが最も影響を受けないかを調べてください)。GPU-off-chip-MEM

コード実行のスレッド化-3D-「ジオメトリ」を変更することはできますが、レイテンシに関して最も重要なリソース ( GPU-SM-REGISTERs ) は制限されており、スケジューリング中のコンテキスト スワップで他のスレッドによって「理想的に再利用」することはできません。計画するスレッドが多いほどGPU-SM-REGISTER、スレッド専用にできる s が少なくなり (それぞれの計算互換性 XY の全体的な静的制限は問題ではありません)、実際のコード中にオフザチップ メモリ アクセスが発生する可能性が高くなります。実行 (SLOCこの事実が与えられるためにa を書く必要はありません。公開されているアーキテクチャ ドキュメントに従ってください)。


A3:いいえ。カーネル分離のアイデアは、別のスレッド 3D ジオメトリ配置からの潜在的な利点の錯覚を導入する可能性がありますが、コードには、データの読み込み/共有コプロセッシング/パブリッシュのために追加のパフォーマンス コストを支払うことで、より多くの問題が発生します。構造。カーネル分離は、完全に設計されたRDMAコード実行で意味があります。

于 2016-03-29T10:30:49.057 に答える