0

CUDA で並列マージ アルゴリズムを実装しようとしています。このアルゴリズムは、1 つのスレッド ブロックで実行されるように設計されています。その基本的な考え方は、2 つの入力シーケンスの各要素のグローバル ランクを計算することです。これら 2 つの入力シーケンスは並べ替えられるため、要素のグローバル ランクは、元のシーケンスのインデックスと、二分探索によって計算されるもう一方のシーケンスのランクを足したものに等しくなります。このようなアルゴリズムを実装するための最良の戦略は、2 つのシーケンスを共有メモリにロードして、グローバル メモリの読み取りを減らすことだと思います。ただし、共有メモリを使用するバージョンと共有メモリを使用しないバージョンの 2 つの実装を比較すると、パフォーマンスの向上は見られません。私は何か間違ったことをしているのだろうか。

ハードウェア: GeForce GTX 285、Linux x86_64。両方の実装で 1024 要素の 2 つのシーケンスをマージする時間は約 0.068672 ミリ秒です。

__global__ void localMerge(int * A, int numA,int * B,int numB,int * C){
extern __shared__ int  temp[]; // shared memory for A and B; 
int tx=threadIdx.x;
int size=blockDim.x;
int *tempA=temp;
int *tempB=temp+numA;

int i,j,k,mid;
    //read sequences into shared memory 
for(i=tx;i<numA;i+=size){
    tempA[i]=A[i];
}
for(i=tx;i<numB;i+=size){
    tempB[i]=B[i];
}
__syncthreads();
    //compute global rank for elements in sequence A
for(i=tx;i<numA;i+=size){
    j=0;
    k=numB-1;
    if(tempA[i]<=tempB[0]){
        C[i]=tempA[i];
    }
    else if(tempA[i]>tempB[numB-1]){
        C[i+numB]=tempA[i];
    }
    else{
        while(j<k-1){
            mid=(j+k)/2;
            if(tempB[mid]<tempA[i]){
                j=mid;
            }
            else{
                k=mid;
            }
        }
        //printf("i=%d,j=%d,C=%d\n",i,j,tempA[i]);
        C[i+j+1]=tempA[i];
    }
}   
    //compute global rank for elements in sequence B
for(i=tx;i<numB;i+=size){
    j=0;
    k=numA-1;
    if(tempB[i]<tempA[0]){
        C[i]=tempB[i];
    }
    else if(tempB[i]>=tempA[numA-1]){
        C[i+numA]=tempB[i];
    }
    else{
        while(j<k-1){
            mid=(j+k)/2;
            if(tempA[mid]<=tempB[i]){
                j=mid;
            }
            else{
                k=mid;
            }
        }
        //printf("i=%d,j=%d,C=%d\n",i,j,tempB[i]);
        C[i+j+1]=tempB[i];
    }
}    
}
4

1 に答える 1

4

メモリ内の両方の入力リストを介した並列のきめ細かいバイナリ検索のコレクションに依存するよりも、 「マージパス」アルゴリズムを適用する方が幸運な場合があり__shared__ます。__shared__このアルゴリズムに再利用が存在するものはキャッシュによってかなりうまくキャプチャできるため、この問題にメモリを使用することはそれほど重要ではありません。

このマージアルゴリズムでは、CTAの各スレッドがkマージ結果の出力を生成する責任があるという考え方です。これには、各スレッドの作業がほぼ均一であり、関連するバイナリ検索がかなり粗いという優れた特性があります。

スレッドiは両方の入力リストを一度k*iに検索して、 th出力要素の各リスト内の位置を見つけます。その場合、ジョブは単純です。各スレッドkは、入力リストの項目を順次マージし、それらを出力の場所にコピーしk*iます。

詳細については、 Thrustの実装を参照してください。

于 2013-01-24T23:09:14.953 に答える