CUDA: __syncthreads() перед операцией общей памяти?

Question

Я в довольно плохой ситуации, не имея возможности использовать отладчик CUDA. Я получаю некоторые странные результаты от использования __syncthreads в приложении с одним общим массивом (deltas). Следующий фрагмент кода выполняется в цикле:

__syncthreads(); //if I comment this out, things get funny 
deltas[lex_index_block] = intensity - mean; 
__syncthreads(); //this line doesnt seem to matter regardless if the first sync is commented out or not 
//after sync: do something with the values of delta written in this threads and other threads of this block 

В основном, у меня есть код с перекрывающихся блоков (необходимые в связи с характером алгоритма). Программа компилируется и запускается, но каким-то образом я получаю систематически неправильные значения в областях вертикального перекрытия. Это меня очень сбивает с толку, поскольку я думал, что правильный способ синхронизации - синхронизация после того, как потоки выполнили мою запись в разделяемую память.

Это целая функция:

//XC without repetitions 
template <int blocksize, int order> 
__global__ void __xc(unsigned short* raw_input_data, int num_frames, int width, int height, 
       float * raw_sofi_data, int block_size, int order_deprecated){ 

//we make a distinction between real pixels and virtual pixels 
//real pixels are pixels that exist in the original data 

//overlap correction: every new block has a margin of 3 threads doing less work (only computing deltas) 
int x_corrected = global_x() - blockIdx.x * 3; 
int y_corrected = global_y() - blockIdx.y * 3; 

//if the thread is responsible for any real pixel 
if (x_corrected < width && y_corrected < height){ 

    //  __shared__ float deltas[blocksize]; 
    __shared__ float deltas[blocksize]; 

    //the outer pixels of a block do not update SOFI values as they do not have sufficient information available 
    //they are used only to compute mean and delta 
    //also, pixels at the global edge have to be thrown away (as there is not sufficient data to interpolate) 
    bool within_inner_block = 
      threadIdx.x > 0 
      && threadIdx.y > 0 
      && threadIdx.x < blockDim.x - 2 
      && threadIdx.y < blockDim.y - 2 
      //global edge 
      && x_corrected > 0 
      && y_corrected > 0 
      && x_corrected < width - 1 
      && y_corrected < height - 1 
      ; 


    //init virtual pixels 
    float virtual_pixels[order * order]; 
    if (within_inner_block){ 
     for (int i = 0; i < order * order; ++i) { 
      virtual_pixels[i] = 0; 
     } 
    } 


    float mean = 0; 
    float intensity; 
    int lex_index_block = threadIdx.x + threadIdx.y * blockDim.x; 



    //main loop 
    for (int frame_idx = 0; frame_idx < num_frames; ++frame_idx) { 

     //shared memory read and computation of mean/delta 
     intensity = raw_input_data[lex_index_3D(x_corrected,y_corrected, frame_idx, width, height)]; 

     __syncthreads(); //if I comment this out, things break 
     deltas[lex_index_block] = intensity - mean; 
     __syncthreads(); //this doesnt seem to matter 

     mean = deltas[lex_index_block]/(float)(frame_idx+1); 

     //if the thread is responsible for correlated pixels, i.e. not at the border of the original frame 
     if (within_inner_block){ 
      //WORKING WITH DELTA STARTS HERE 
      virtual_pixels[0] += deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y + 1, 
         blockDim.x)] 
        * 
        deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y - 1, 
         blockDim.x)]; 

      virtual_pixels[1] += deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y, 
         blockDim.x)] 
        * 
        deltas[lex_index_2D(
         threadIdx.x + 1, 
         threadIdx.y, 
         blockDim.x)]; 

      virtual_pixels[2] += deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y, 
         blockDim.x)] 
        * 
        deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y + 1, 
         blockDim.x)]; 

      virtual_pixels[3] += deltas[lex_index_2D(
         threadIdx.x, 
         threadIdx.y, 
         blockDim.x)] 
        * 
        deltas[lex_index_2D(
         threadIdx.x+1, 
         threadIdx.y+1, 
         blockDim.x)]; 
      //    xc_update<order>(virtual_pixels, delta2, mean); 
     } 
    } 

    if (within_inner_block){ 
     for (int virtual_idx = 0; virtual_idx < order*order; ++virtual_idx) { 
      raw_sofi_data[lex_index_2D(x_corrected*order + virtual_idx % order, 
             y_corrected*order + (int)floorf(virtual_idx/order), 
             width*order)]=virtual_pixels[virtual_idx]; 
     } 
    } 
} 
} 

Answer 1

Из того, что я могу видеть, там может быть опасность в вашем приложении между итераций цикла. Запись в deltas[lex_index_block] для циклической итерации frame_idx+1 может быть сопоставлена ​​с тем же местом, что и для чтения deltas[lex_index_2D(threadIdx.x, threadIdx.y -1, blockDim.x)] в другом потоке на итерации frame_idx. Два обращения неупорядочены, и результат является недетерминированным. Попробуйте запустить приложение с cuda-memcheck --tool racecheck.