PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

Warning: Cannot tell what pointer points to

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
TheMatt



Joined: 06 Jul 2009
Posts: 322
Location: Greenbelt, MD

PostPosted: Mon Nov 08, 2010 11:58 am    Post subject: Warning: Cannot tell what pointer points to Reply with quote

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:
Code:
./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:
Code:
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:
Code:
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
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6215
Location: The Portland Group Inc.

PostPosted: Tue Nov 09, 2010 3:22 pm    Post subject: Reply with quote

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
Back to top
View user's profile
TheMatt



Joined: 06 Jul 2009
Posts: 322
Location: Greenbelt, MD

PostPosted: Wed Nov 10, 2010 5:47 am    Post subject: Reply with quote

mkcolg wrote:
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).

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.)

Quote:
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.
Back to top
View user's profile
TheMatt



Joined: 06 Jul 2009
Posts: 322
Location: Greenbelt, MD

PostPosted: Wed Nov 10, 2010 11:39 am    Post subject: Reply with quote

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:
Code:
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:
Code:
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
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming All times are GMT - 7 Hours
Page 1 of 1

 
Jump to:  
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum


Powered by phpBB © phpBB Group