vendredi 8 mai 2015

CUDA - nvcc -G - if not working properly

I'm currently working on porting a lava flow model in CUDA (full code on github here: Full source of the CUDA-SCIARA Fv2 lava flow model . I'm facing a problem with an if construct inside a __device__ function.

Based on the current lava quantity and temperature it computes the new temperature and if it is lower than a constant parameter (the variable d_PTsol=1143.0) the lava is solified.

The problem in the code below is that is works perfectly if I compile with the -G options (for the generation of device code debug infos) but behave wrongly without.

double new_temp = d_computeNewTemperature(sommah,sommath);        
if(new_temp <= d_PTsol){
            printf("Solidified %.5f,%.5f\n",new_temp,d_PTsol);
            double newQuote = d_sbts_updated[d_getIdx(row,col,ALTITUDE)]+d_sbts_current[d_getIdx(row,col,THICKNESS)];
            //CODE FOR LAVA SOLIDIFICATION HERE
    }else{
           //there is lava and is not solidified -> activate this cell!
           adjustAdaptiveGrid(row,col);
 }

ouptutting something like this at a certain point of the simulation:

Solidified 1344.68654 1143.00000
Solidified 1343.99509 1143.00000
Solidified 1320.50061 1143.00000
Solidified 1325.53942 1143.00000

To make things more subtle the problem completly disappear if I change the the if condition to a strict inequality if(new_temp < d_PTsol).

Compilation is carried out with the following options and in separate compilation mode

-O3 -Xcompiler -fPIC -std=c++11

and linking using

--cudart static --relocatable-device-code=true -gencode arch=compute_35,code=compute_35 -gencode arch=compute_35,code=sm_35

Has someone faced a similar issue before? Am I doing something wrong?

UPDATE

The problem seem to be somehow related to the translation of the if else construct with the <= as condition. Translating

if(new_temp <= d_PTsol) {
        //solidification
}else{
        //something else
}

to

if(new_temp <= d_PTsol) {
        //solidification
}
if(!(new_temp <= d_PTsol)){
        //something else
}

makes the code work perfectly.

Aucun commentaire:

Enregistrer un commentaire