i'm working on porting lava flow model in cuda (full code on github here: full source of cuda-sciara fv2 lava flow model .
edit
in order reproduce issue, use instruction on github readme in order dev copy of project sciara_fv2_cuda_multicells. compile withouth -g option , passing command line argument -c ../data/2006/parameters.cfg . (quoted code @ line 260 in file )
i'm facing problem if construct inside __device__
function.
based on current lava quantity , temperature computes new temperature , if lower constant parameter (the variable d_ptsol=1143.0
) lava solified.
the problem in code below is works if compile -g options (for generation of device code debug infos) 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 lava solidification here }else{ //there lava , not solidified -> activate cell! adjustadaptivegrid(row,col); }
ouptutting @ point of 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 problem completly disappear if change the if condition strict inequality if(new_temp < d_ptsol)
.
compilation carried out following options , 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 faced similar issue before? doing wrong?
update
the problem seem somehow related translation of if else construct <= 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 code work perfectly.
i had similar issue. code worked -g option not without. needed add -fmad=false , -prec-div=false compiler options (for more information see: http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#options-for-steering-gpu-code-generation).
by default values true , optimise single precision operations, @ cost of precision. if enable debug mode optimisation disabled. needed exact values, therefore code worked -g. maybe have similar issue. (more information cuda floating point precision: http://docs.nvidia.com/cuda/cuda-c-programming-guide/#mathematical-functions-appendix)
Comments
Post a Comment