Warning: Cannot tell what pointer points to

I’m currently trying to move a code to GPUs that, like all my code, is quite complex. With this latest attempt, I am seeing a new Warning on the command line that I’ve never seen before:

./cloudnew.001.gpu(1547): Warning: Cannot tell what pointer points to, assuming global memory space
./cloudnew.001.gpu(1548): Warning: Cannot tell what pointer points to, assuming global memory space
./cloudnew.001.gpu(1551): Warning: Cannot tell what pointer points to, assuming global memory space
./cloudnew.001.gpu(1556): Warning: Cannot tell what pointer points to, assuming global memory space
./cloudnew.001.gpu(1563): Warning: Cannot tell what pointer points to, assuming global memory space

I’m a bit puzzled by these. The relevant code in the .gpu file is:

1537 extern "C" __device__ float dqsat(signed char* _ptl,signed char* _ppl,signed char* _pqsat){
1538 float qi;
1539 float dqi;
1540 float dqq;
1541 float qq;
1542 float _rdqsat;
1543 float ti;
1544 float pp;
1545 float uramp;
1546 uramp = _cloudnew_17.m9908;
1547 pp = __fmul_rn((*(float*)(_ppl)), (1.00000000e+02f));
1548 ti = (*(float*)(_ptl))-_cloudnew_17.m9880;
1549 if( !((ti)<=(uramp))) goto _BB_5;
1550 qq = qsatice0(_ptl, ((signed char*)(&pp)), ((signed char*)(&dqsat)));
1551 *(float*)(_pqsat) = qq;
1552 goto _BB_9;
1553 _BB_5: ;
1554 if( !((ti)>=((0.00000000e+00f)))) goto _BB_7;
1555 qq = qsatlqu0(_ptl, ((signed char*)(&pp)), ((signed char*)(&dqsat)));
1556 *(float*)(_pqsat) = qq;
1557 goto _BB_9;
1558 _BB_7: ;
1559 qq = qsatlqu0(_ptl, ((signed char*)(&pp)), ((signed char*)(&dqq)));
1560 qi = qsatice0(_ptl, ((signed char*)(&pp)), ((signed char*)(&dqi)));
1561 ti = ti/uramp;
1562 _rdqsat = dqq+(__fmul_rn(ti, dqi-dqq));
1563 *(float*)(_pqsat) = qq+(__fmul_rn(ti, qi-qq));
1564 _BB_9: ;
1565 return _rdqsat;
1566 }

and the “real” code is:

3277    attributes(device) function DQSAT(TL,PL,QSAT)
3278       real,   intent(IN) :: TL, PL
3279       real,   intent(OUT):: QSAT
3280       real    :: DQSAT
3281       real    :: URAMP, TT, WW, DD, DQQ, QQ, TI, DQI, QI, PP
3282       integer :: IT
3283 
3284       URAMP = TMIX
3285 
3286       PP = PL*100.
3287 
3288       TI = TL - ZEROC
3289 
3290       if    (TI <= URAMP) then
3291          QQ  = QSATICE0(TL,PP,DQ=DQSAT)
3292          QSAT  = QQ
3293       elseif(TI >= 0.0  ) then
3294          QQ  = QSATLQU0(TL,PP,DQ=DQSAT)
3295          QSAT  = QQ
3296       else
3297          QQ  = QSATLQU0(TL,PP,DQ=DQQ)
3298          QI  = QSATICE0(TL,PP,DQ=DQI)
3299          TI  = TI/URAMP
3300          DQSAT = TI*(DQI - DQQ) + DQQ
3301          QSAT  = TI*(QI - QQ) +  QQ
3302       end if
3303 
3304    end function DQSAT

So, the warning is thrown when one of the inputs is mentioned, it seems. Now, I will say that there is a QSAT function just above this that looks kinda similar…could I be seeing a name clash? Or maybe it’s related to the other warning I see with this compile?

Any help is appreciated,
Matt

Hi Matt,

It looks like the NVIDIA compiler can’t tell which memory space the TL, PL, and QSAT are located in. This is an issue on Telsa’s since it has distinct instructions for loads from global memory, constant memory, and shared memory. Fermi has a unified memory model so probably wouldn’t get the warning.

You can fix this for TL and PL by passing them in by value instead of using intent(in).

If QSAT is a ‘shared’ variable, then you might have some issues and need to rework the code.

  • Mat

I’ll test this out, thanks.

Is this something I should do/need to do for all the device functions I have? And does that include subroutines as well? I ask because with this code I managed to completely scalarize the internals so all the device subs and functions are now all scalar inputs and outputs.

ETA: This does indeed help out with the Warnings regarding TL and PL, though I’m not sure if these will always be in global memory–as I thought values had to be–as they could be registers. (I’m fairly certain I’m under register pressure, so it may be true these get pushed to global.)

If QSAT is a ‘shared’ variable, then you might have some issues and need to rework the code.

It isn’t shared (yet) so that’s not the issue. Hmm. I’m sending in some code in regards to my other thread, so maybe you can take a look and see what you think.

Problem possibly solved (or at least, warnings not there anymore).

The issue seems to have been due to the use of named keyword arguments (or whatever they are called) in calling these functions.

For example, in my code I had this:

DQSx  = DQSAT( TEo, PL, QSAT=QSx )

which is correct for the CPU because the DQSAT it calls has many optional arguments (as well as being overloaded for 0-, 1-, 2-, and 3-D input/outputs). For the GPU I had stripped this functionality out and kept only the bare code needed for DQSAT to replicate what was going on. Likewise, in the DQSAT code I had:

QQ  = QSATICE0(TL,PP,DQ=DQSAT)

where the original QSATICE0 also had optional arguments. I imagine I made mistakes in this, or the compiler was just getting confused somehow.

So, I went through all the code here and got rid of all this X=Y jazz and just did “normal” function references. Doing this meant no more warnings thrown and (promisingly?) the compile time for the GPU is now c. 2 minutes instead of c. 15s.

I’m still not sure the code works, but it compiles differently at least.

Thanks,
Matt