CUDA __syncthreads(); 작동 안함; 역방향 중단 점 히트 주문

Question

나는 __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 거의 변화, http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html에서입니다.

이제는 3 개의 중단 점을 설정했습니다. A, B, C. 그것은 A, B, C 순으로 공격해야합니다. 문제는 A, B * x, C, B 순으로 맞았다는 것입니다. 따라서 C 지점에서 selectedRecordsOffset은 준비되지 않았으므로 오류가 발생합니다. A 후에 B는 몇 번 누르지 만 전부는 C가 안되면 코드에서 더 나아가고 나머지 B에서는 다시 B가됩니다. x는 입력에 따라 다릅니다 (일부 입력에 대해서는 중단 점에 역수가 없으므로 C가 마지막에 도달했습니다).

더욱이 나는 A와 C의 threadIdx.y = 0과 B threadIdx.y = 10에 대해 히트를 일으키는 스레드 번호를 살펴 본다. 이것이 동일한 블록 인 동안 어떻게 가능 할까? 그래서 일부 스레드는 동기화되지 않는 이유는 무엇인가? ? 조건부 동기화는 없습니다. 누군가 버그를 찾을 수있는 아이디어가 있습니까?

더 자세한 설명이 필요하면 질문하십시오.
미리 조언 해 주셔서 감사합니다.
아담

Answer 1

Thou shalt not use __syncthreads() in conditional code 각 블록의 모든 스레드에서 균일하게 평가되지 않는 경우.