2分木を累積的に構築する

Question

ベクトルs.tで2分木を構築したい。親の値は両方の子の合計になります。再帰的にCでツリーを構築するには、次のようになります。

int construct(int elements[], int start, int end, int* tree, int index) { 
    if (start == end) { 
     tree[index] = elements[start]; 
     return tree[index]; 
    } 
    int middle = start + (end - start)/2; 
    tree[index] = construct(elements, start, middle, tree, index*2) + 
        construct(elements, middle, end, tree, index*2+1); 
    return tree[index]; 
} 

しかし、私はスレッドを利用することにより、並列の方法でCUDAでそれを構築する方法を知りません。 1 reference有用であるとわかりました

この種の再帰アルゴリズムをどのように並列化するべきですか？ 1つの方法は、ガランザ（Garanzha）らによって提示された方法を使用することであり、これはルートから順にノードのレベルを処理する。この考え方は、幅広いノードの配列を幅優先順に維持することで、階層内のすべてのレベルが線形範囲のノードに対応するようにします。与えられたレベルでは、この範囲に入る各ノードごとに1つのスレッドを起動します。スレッドは、最初と最後をノード配列から読み込み、findSplit（）を呼び出すことで開始します。次に、結果の子ノードをアトミ​​ックカウンタを使用して同じノード配列に追加し、対応するサブ範囲を書き出します。このプロセスは、各レベルが次のラウンドで処理される次のレベルに含まれるノードを出力するように反復する。

各レベルを順次処理し、各レベルでノードを並列化します。私はそれが完全に意味をなさないと思うが、私はそれを正確に実装する方法はありません、誰かが私にそれを行う方法のアイデアや例を教えてもらえますか？

Answer 1

上記のインデックススキームが機能するかどうかはわかりません。ここで

は仕事ができるサンプルコードです：（ツリー索引がニーズに合わないかもしれませんが）：

__global__ void buildtreelevel(const int* elements, int count, int* tree) 
{ 
    int parentcount = (count + 1) >> 1; 
    for (int k = threadIdx.x + blockDim.x * blockIdx.x ; k < parentcount ; k += blockDim.x * gridDim.x) 
    { 
     if ((2*k+1) < count) 
      tree[k] = elements[k*2] + elements[k*2+1] ; 
     else 
      tree[k] = elements[k*2] ; 
    } 
} 

この機能は、一度に1つのツリーレベルを処理します。全体的な木の大きさは、によって提供されています

int treesize (int count, int& maxlevel) 
{ 
    int res = 1 ; 
    while (count > 1) 
    { 
     count = (count + 1) >> 1 ; 
     res += count ; 
     ++maxlevel; 
    } 
    return res ; 
} 

とツリー全体を構築することがbuildtreelevelカーネルには、いくつかの呼び出しを必要とします。

int buildtree (int grid, int block, const int* d_elements, int count, int** h_tree, int* d_data) 
{ 
    const int* ptr_elements = d_elements ; 
    int* ptr_data = d_data ; 
    int level = 0 ; 
    int levelcount = count ; 
    while (levelcount > 1) 
    { 
     buildtreelevel <<< grid, block >>> (ptr_elements, levelcount, ptr_data) ; 
     levelcount = (levelcount + 1) >> 1 ; 

     h_tree [level++] = ptr_data ; 
     ptr_elements = ptr_data ; 
     ptr_data += levelcount ; 
    } 
    return level ; 
} 

同期は唯一、すべてのカーネルがストリーム上で実行される最後で発生する必要があります0

int main() 
{ 
    int nElements = 10000000 ; 
    int* d_elements ; 
    int* d_data ; 
    int** h_tree ; 
    int maxlevel = 1 ; 

    cudaMalloc ((void**)&d_elements, nElements * sizeof (int)) ; 
    cudaMalloc ((void**)&d_data, treesize(nElements, maxlevel) * sizeof (int)) ; 

    h_tree = new int*[maxlevel]; 

    buildtree (64, 256, d_elements, nElements, h_tree, d_data) ; 

    cudaError_t res = cudaDeviceSynchronize() ; 
    if (cudaSuccess != res) 
     fprintf (stderr, "ERROR (%d) : %s \n", res, cudaGetErrorString(res)); 
    cudaDeviceReset(); 
} 

ツリー構造は、デバイスポインタのホスト配列であるh_treeに格納されます。

これは最適ではありませんが、__ldgでアライメントされたint4を使用することをお勧めします。一度に4レベルを処理すると、パフォーマンスが向上する可能性があります。