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，几乎没有变化。

好的，现在我设置了 3 个断点：A、B、C。它们应该按 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 回答