1

プレフィックスサムを計算するためのコードを書いています。ここに私のカーネルがあります 

__global__ void prescan(int *indata,int *outdata,int n,long int *sums)  
{  
    extern __shared__ int temp[];  

    int tid=threadIdx.x;
    int offset=1,start_id,end_id;
    int *global_sum=&temp[n+2];

    if(tid==0)
    {
        temp[n]=blockDim.x*blockIdx.x;
        temp[n+1]=blockDim.x*(blockIdx.x+1)-1;
        start_id=temp[n];
        end_id=temp[n+1];
        //cuPrintf("Value of start %d and end %d\n",start_id,end_id);

    }
    __syncthreads();
    start_id=temp[n];
    end_id=temp[n+1];
    temp[tid]=indata[start_id+tid];
    temp[tid+1]=indata[start_id+tid+1];


    for(int d=n>>1;d>0;d>>=1)
    {
        __syncthreads();
        if(tid<d)
        {
            int ai=offset*(2*tid+1)-1;
            int bi=offset*(2*tid+2)-1;

            temp[bi]+=temp[ai];
        }
        offset*=2;
    }

    if(tid==0)
    {  
        sums[blockIdx.x]=temp[n-1];  
        temp[n-1]=0;
        cuPrintf("sums %d\n",sums[blockIdx.x]);
    }
    for(int d=1;d<n;d*=2)
    {
        offset>>=1;
        __syncthreads();
        if(tid<d)
        {
            int ai=offset*(2*tid+1)-1;
            int bi=offset*(2*tid+2)-1;
            int t=temp[ai];
            temp[ai]=temp[bi];
            temp[bi]+=t;
        }
    }

    __syncthreads();

    if(tid==0)
    {
        outdata[start_id]=0;
    }

    __threadfence_block();
    __syncthreads();
    outdata[start_id+tid]=temp[tid];
    outdata[start_id+tid+1]=temp[tid+1];
    __syncthreads();

    if(tid==0)  
    {  
        temp[0]=0;  
        outdata[start_id]=0;  

    }  

    __threadfence_block();
    __syncthreads();

    if(blockIdx.x==0 && threadIdx.x==0)
    {
        for(int i=1;i<gridDim.x;i++)
        {
            sums[i]=sums[i]+sums[i-1];
        }
    }

    __syncthreads();
    __threadfence();

    if(blockIdx.x==0 && threadIdx.x==0)
    {
        for(int i=0;i<gridDim.x;i++)
        {
            cuPrintf("****sums[%d]=%d ",i,sums[i]);
        }
    }


    __syncthreads();
    __threadfence();


    if(blockIdx.x!=gridDim.x-1)
    {
        int tid=(blockIdx.x+1)*blockDim.x+threadIdx.x;
        if(threadIdx.x==0)
            cuPrintf("Adding %d \n",sums[blockIdx.x]);
        outdata[tid]+=sums[blockIdx.x];

    }
    __syncthreads();

}

上記のカーネルでは、合計配列はブロックごとにプレフィックス合計を累積し、次に最初のスレッドがこの合計配列のプレフィックス合計を計算します。デバイス側からこの合計配列を出力すると、正しい結果が表示されます

cuPrintf("%d \n を追加中",sums[blockIdx.x]);

この行は、古い値を取得していることを出力します。その理由は何ですか？

4

1 に答える 1

2

あなたのコードは、マルチブロック プレフィックス サムの有効な実装ではありません。ブロック 0 の単一スレッドを使用して、ブロックの部分合計がメモリに書き込まれたことが保証される前に、ブロックの合計のプレフィックス合計を計算しようとしています。 __syncthreads()ブロック間ではなく、単一のブロック内のスレッドのみを同期します。したがって、このコードでは：

__threadfence_block();
__syncthreads();

if(blockIdx.x==0 && threadIdx.x==0)
{
    for(int i=1;i<gridDim.x;i++)
    {
        sums[i]=sums[i]+sums[i-1];
    }
}

ブロック 0 がこのコードを実行する前に、すべてのブロックが合計 [blockIdx.x] を計算したとは限りません。実際、デバイスで同時に実行できるよりも多くのブロックを起動すると、このコードに到達したときにすべてのブロックが開始されているとは限りません。

このコードを正しくするには、このコードの前にカーネルを終了し、新しいカーネルを起動してブロック プレフィックスの合計結果を計算し、別のカーネルを起動してその結果を各スレッド ブロックに追加する必要があります。

于 2012-08-31T03:35:47.690 に答える