PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

copyin copyout with CUDA Fortran

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling
View previous topic :: View next topic  
Author Message
Tuan



Joined: 11 Jun 2009
Posts: 233

PostPosted: Sun Feb 09, 2014 2:29 pm    Post subject: copyin copyout with CUDA Fortran Reply with quote

Recently, when I turn on the flag -Minfo, I see a lot of 'copyin' and 'copyout' for each functional call.

Code:
setup_ryr_2:                                                                               
   2825, Possible copy in and copy out of akrm in call to getcompk_4                                                                                         
   2826, Possible copy in and copy out of akrp in call to getcompk_4   

put_sfu2grid3d_nonuniform_rogue_ryr:                                                   
   2246, Copy in and copy out of sfu_coords in call to valid_sfu_location                                                                                 
   2247, Copy in and copy out of sfu_coords in call to map_loc2grid                                                                                       
   2283, Copy in and copy out of sfu_coords in call to valid_sfu_location                                                                                 
   2284, Copy in and copy out of sfu_coords in call to map_loc2grid           



I'm not using Accelerator programming model, but CUDA Fortran. So, my question is how can I optimize the code, written in CUDA Fortran, to specify the compiler when to do copyin and/or copyout?
Can it be done using clauses like in Fortran Accelerator? Please give me an example if it can be done.

Thanks,
Tuan
Back to top
View user's profile
mkcolg



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

PostPosted: Tue Feb 11, 2014 4:53 pm    Post subject: Reply with quote

Hi Tuan,

These have nothing to do with CUDA Fortran or OpenACC. The compiler is tell you that it may in one case and has in another created a temp array to pass in a non-unit 1 array (non-contiguous) array to a subroutine. This can cause performance problems which can only be fixed by passing in contiguous arrays.

- Mat
Back to top
View user's profile
Tuan



Joined: 11 Jun 2009
Posts: 233

PostPosted: Fri Feb 21, 2014 4:54 pm    Post subject: Reply with quote

Hi Mat,
Based on your response I have a question. Padding is quite often being used in C. So, my question is if I use padding and pass an device array like this in Fortran,

Code:

real, device :: matA

allocate(matA(pad+size))

call my_kernel <<<...,...>>> (matA(pad+1:), ...)



should the CUDA Fortran compiler create a duplicate of matA(pad+1:*) or it just use the original matA array. Can the compiler detect that matA(pad+1:*) is still a contiguous array here?

Thanks,
Tuan
Back to top
View user's profile
mkcolg



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

PostPosted: Mon Feb 24, 2014 1:42 pm    Post subject: Reply with quote

Hi Tuan,

Quote:
Can the compiler detect that matA(pad+1:*) is still a contiguous array here?
Yes, however, the compiler may still have to create a new descriptor if your interface for this array is assumed-size (*).

Quote:
should the CUDA Fortran compiler create a duplicate of matA(pad+1:*) or it just use the original matA array.
In CUDA Fortran, it actually can't duplicate the array. However, since the contiguous test can be done until run time, you may get run time failures if a reshape is needed or the test can determine if it's contiguous.

- Mat
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling 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