Problem with sum of two structs without ripetition

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi people, the following code does the sum of two structs without ripetition and, for each couple, it does an “if” that, if success, stores the result in a vector:

__device__ void createPI0(QVECT *Gamma, PI_0 *Pi0, unsigned long int N, unsigned long int N2, float M0, float del, int *dimPi0){

//Dichiarazioni variabili locali sul device

        int linIdx, i, z, j, loc, sh, tid;

        float massa;

        QVECT Candidato;

//Dichiarazione variabile condivisa sul device

        __shared__ int Ind;

//Il primo thread del blocco inizializza la variabile shared

        if (threadIdx.x == 0) {

           Ind = *dimPi0;

        }

        __syncthreads(); //Barriera di sincronizzazione threads

//Calcolo del thread ID

        sh=threadIdx.x+blockDim.x*threadIdx.y;

        tid = blockDim.x*blockDim.y*(blockIdx.x + gridDim.x* blockIdx.y) + sh;

//Algoritmo per il calcolo degli indici giusti in base al thread

        if (tid<N2){

                linIdx=N2-tid;

                i=int(N - 0.5 - sqrt(0.25 - 2 * (1 - linIdx)));

                z=(N+N-1-i)*i;

                j=tid - z/2 + 1 + i;

if (i==j){

                        i=i-1;

                        j=N-1;

                }

//Somma di due quadrivettori

                Candidato.x=Gamma[i].x+Gamma[j].x;

                Candidato.y=Gamma[i].y+Gamma[j].y;

                Candidato.z=Gamma[i].z+Gamma[j].z;

                Candidato.Ene=Gamma[i].Ene+Gamma[j].Ene;

//Controllo massa

                massa=(Candidato.Ene*Candidato.Ene)-(Candidato.x*Candidato.x)-(Candidato.y*Candidato.y)-(Candidato.z*Candidato.z);

if(massa>M0-del && massa<M0+del){

                        loc=atomicAdd(&Ind,1);

                        Pi0[loc].x=Candidato.x;

                        Pi0[loc].y=Candidato.y;

                        Pi0[loc].z=Candidato.z;

                        Pi0[loc].Ene=Candidato.Ene;

                        Pi0[loc].g1=i;

                        Pi0[loc].g2=j;

                }

        }

__syncthreads();

        //Il thread 0 di ogni blocco aggiorna la dimensione parziale dei Pi 0

        if(threadIdx.x==0)

                atomicAdd(dimPi0,Ind);

}

Now the trouble is that even 10000 threads works fine, for more of 10000 threads the results aren’t correct. Each thread manages only a sum, only a couple of element. Thanks for the help!

2

no solution? :(