2011-10-31 10 views
2

ここに私の問題があります:私はcudaのどこかに格納されるかなり大きな2倍(77.500倍の配列)です。今度は、その配列で一連の操作を順番に実行するために、大きなスレッドセットが必要です。すべてのスレッドは、その配列のSAME要素を読み取り、タスクを実行し、結果を共有メモリに格納し、配列の次の要素を読み取る必要があります。すべてのスレッドは、同じメモリ位置から同時に読み取る(ちょうど読み取る)必要があることに注意してください。だから私は疑問に思う:ただ1つのメモリを読み込んで同じダブルをすべてのスレッドにブロードキャストする方法はありますか?何度も何度も読むのは無用だと思います。CUDA:すべてのスレッドに同じメモリ位置をストリーミングする

答えて

3

これは一般的な最適化です。アイデアは、各スレッドがデータを読み込むために、そのblockmatesと協力することです:

// choose some reasonable block size 
const unsigned int block_size = 256; 

__global__ void kernel(double *ptr) 
{ 
    __shared__ double window[block_size]; 

    // cooperate with my block to load block_size elements 
    window[threadIdx.x] = ptr[threadIdx.x]; 

    // wait until the window is full 
    __syncthreads(); 

    // operate on the data 
    ... 
} 

あなたが「スライド」(以上または多分いくつかの整数倍)配列block_size間でウィンドウを反復することができます消費する時の要素全部。同じテクニックは、データを同期して保存する場合にも適用されます。

+2

また、共有メモリからのブロードキャスト(つまり、ブロック内のすべてのスレッドが同じメモリ位置を読み取る)は高速です。 N身体問題のCUDA実装では、Jaredがここで説明したイディオムと一緒にブロードキャストを使用します。 – ArchaeaSoftware

関連する問題