0

__syncthreads();次のコードに問題があると思います:

__device__ void prefixSumJoin(const bool *g_idata, int *g_odata, int n)
{
    __shared__ int temp[Config::bfr*Config::bfr];  // allocated on invocation  
    int thid = threadIdx.y*blockDim.x + threadIdx.x;  
    if(thid<(n>>1))
    {
        int offset = 1;
        temp[2*thid] = (g_idata[2*thid]?1:0); // load input into shared memory  
        temp[2*thid+1] = (g_idata[2*thid+1]?1:0); 
        for (int d = n>>1; d > 0; d >>= 1)                    // build sum in place up the tree  
        {   
            __syncthreads();  
            if (thid < d)  
            { 
                int ai = offset*(2*thid+1)-1; // <-- breakpoint B 
                int bi = offset*(2*thid+2)-1;
                temp[bi] += temp[ai];  
            }  
            offset *= 2; 
        } 
        if (thid == 0) { temp[n - 1] = 0; } // clear the last element

        for (int d = 1; d < n; d *= 2) // traverse down tree & build scan  
        {  
            offset >>= 1;  
            __syncthreads();  
            if (thid < d)                       
            {
                int ai = offset*(2*thid+1)-1;  
                int bi = offset*(2*thid+2)-1;
                int t = temp[ai];  
                temp[ai] = temp[bi];  
                temp[bi] += t;   
            }  
        }  
        __syncthreads();
        g_odata[2*thid] = temp[2*thid]; // write results to device memory  
        g_odata[2*thid+1] = temp[2*thid+1]; 
    }
}


__global__ void selectKernel3(...)
{
    int tidx = threadIdx.x;
    int tidy = threadIdx.y;
    int bidx = blockIdx.x;
    int bidy = blockIdx.y;
    int tid = tidy*blockDim.x + tidx;
    int bid = bidy*gridDim.x+bidx;
    int noOfRows1 = ...;
    int noOfRows2 = ...;

    __shared__ bool isRecordSelected[Config::bfr*Config::bfr];
    __shared__ int selectedRecordsOffset[Config::bfr*Config::bfr];

    isRecordSelected[tid] = false;
    selectedRecordsOffset[tid] = 0;


    __syncthreads();
    if(tidx<noOfRows1 && tidy<noOfRows2)
        if(... == ...)
            isRecordSelected[tid] = true;

    __syncthreads();
    prefixSumJoin(isRecordSelected,selectedRecordsOffset,Config::bfr*Config::bfr); // <-- breakpoint A
    __syncthreads();

    if(isRecordSelected[tid]==true){
        {
            some_instruction;// <-- breakpoint C
        ...
        }
    }
}
...
f(){
   dim3 dimGrid(13, 5);
   dim3 dimBlock(Config::bfr, Config::bfr);

   selectKernel3<<<dimGrid, dimBlock>>>(...)

}
//other file

class Config
{
public:
    static const int bfr = 16; // blocking factor = number of rows per block
public:
    Config(void);
    ~Config(void);
};

prefixSum は GPU Gems 3: Parallel Prefix Sum (Scan) with CUDAからのもので、ほとんど変更されていません。

では、A、B、C の 3 つのブレークポイントを設定しました。A、B、C の順序でヒットする必要があります。問題は、A、B*x、C、B の順序でヒットすることです。ポイント C、selectedRecordsOffset の準備ができておらず、エラーが発生します。A の後、B が数回ヒットしますが、すべてではありません。次に C がヒットし、コード内でさらに進み、ループの残りの部分で B が再びヒットします。x は入力によって異なります (一部の入力では、ブレークポイントに反転がないため、C が最後にヒットしました)。

さらに、ヒットを引き起こすスレッド番号を見ると、A と C threadIdx.y= 0 と B threadIdx.y= 10 です。同じブロックであるのに、なぜ一部のスレッドが同期を省略しているのでしょうか? 条件付き同期はありません。バグを探す場所について何か考えがある人はいますか?

さらに明確にする必要がある場合は、質問してください。
これを解決する方法についてのアドバイスを事前に感謝します。
アダム

4

1 に答える 1

4

条件が各ブロックのすべてのスレッドで均一に評価されない場合は、条件付きコードで __syncthreads() を使用しないでください。

于 2013-05-30T00:59:12.090 に答える