Различные результаты каждого запуска в CUDA

Question

Я программирую в cuda следующее:

1. Я получаю матрицу со значениями изображения (d_Data)
2. Копирую матрицу в разделяемую память (тайл)
3. Я получаю базовую разницу в пикселях (pixel(i)-pixel(i+1)) (d_diff)
4. Если разность является конкретным значением для примера 0, число 1 устанавливается в матрице (d_A) в каждой позиции, где число 0 было найдено в d_diff. Это необходимо для получения частоты 0 в исходной матрице различий.
5. Параллельная накопленная сумма.
6. Результат частоты перейти к вектору частот.

Проверяя шаг за шагом, все идет как положено до накопленной суммы. Когда я запускаю код, программно рассчитанное значение составляет 104347, но иногда из CUDA я получаю результат nan, в других случаях я получаю любое число, например, 2425. Очень странно, если я продолжаю работать с ядром 20 или 30 раз, значение становится ожидаемое 104347: с.

Я использую для каждой матрицы:

h_Data  = (float *)malloc(data_size);
h_diff  = (float *)malloc(data_size);
h_A         = (float *)malloc(data_size);

а также

 cudaFree(d_A);
cudaFree(d_diff);
cudaFree(d_Av);

поэтому я не понимаю, почему код все ближе и ближе к правильному результату, когда я запускаю достаточное количество раз. Кстати, когда достигается правильное значение, оно больше не движется, независимо от того, сколько раз я запускаю код.

Код:

 __global__ void spam(float *d_Data, float *d_diff, float *d_A, int dw, int dh, float *d_Av){

long bx = blockIdx.x;  long by = blockIdx.y;
long tx = threadIdx.x; long ty = threadIdx.y;


// Identify the row and column of the Pd element to work on
long Row = by * TILE_WIDTH + ty;
long Col = bx * TILE_WIDTH + tx;
long tid = Row*dw+Col;
long i=512*512;
long r = MASK_DIM/2;
long s = 0;

 __shared__ int tile[BLOCK_WIDTH][BLOCK_WIDTH];

for (int k=0; k<=8; k++)
     d_Av[k]=0; 


    if(tid < dw*dh)
    {

   // to shared memory.
                                          tile[ty + r][tx + r]=d_Data[Row*dw+Col];
        if (Col-r >=0)                    tile[ty + r]  [tx] = d_Data[Row*dw+Col-r];
        if (Col+r <dw)                    tile[ty + r]  [tx + 2*r] = d_Data[Row*dw+Col+r];
        if (Row-r >=0)                    tile[ty]      [tx + r] = d_Data[(Row - r)*dw + Col];
        if (Row+r <dw)                    tile[ty + 2*r][tx + r] = d_Data[(Row + r)*dw + Col];
        if (Row - r >= 0 && Col - r >= 0) tile[ty]      [tx] = d_Data[(Row-r)*dw+Col-r];
        if(Row - r >= 0 && Col + r < dw)  tile[ty]      [tx + 2*r] = d_Data[(Row-r)*dw+Col+r];
        if (Row + r < dw && Col - r >= 0) tile[ty + 2*r][tx] = d_Data[(Row+r)*dw+Col-r];
        if(Row + r <dw && Col + r < dw)   tile[ty + 2*r][tx + 2*r] = d_Data[(Row-r)*dw+Col+r]; 

        //Calculates the difference matrix
       d_diff[tid] = (tile[ty + r][tx +r] - tile[ty + r][tx + r + 1]);


        d_A[tid]=0;

       //Set a 1 in each position in d_A where 0 was found in d_diff.
        if (d_diff[tid] == 0)
        { d_A[tid]=1;}
        __syncthreads();

        //cumulative sum to get the frecuency of value 0 in d_diff.  // The error is HERE
      for (s = (i/2); s>=1; s=s/2) {
            if (tid < s)
            {   d_A[tid] += d_A[tid+s];
            }
        } 

       // set the frequency value in frequencies vector.
        d_Av[0] = d_A[0];

}} // END IF tid < dw*dh

Любая идея приветствуется:D

Answer 1

Вы можете попробовать заменить оператор if следующим кодом:d_A[tid] += d_A[tid+s] * (tid < s);

И будьте уверены, что этот код не приводит к состоянию гонки. Часто это может быть случай с параллельной суммой.

MK