PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Using Shared Memory

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



Joined: 07 Apr 2012
Posts: 29

PostPosted: Tue Jun 12, 2012 12:07 pm    Post subject: Using Shared Memory Reply with quote

Hello,

I want to use shared memory in acc region, it would be nice if I have an example. I think there is no "shared" subclause, like "private".

Thank you,
Sayan
Back to top
View user's profile
mkcolg



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

PostPosted: Tue Jun 12, 2012 12:53 pm    Post subject: Reply with quote

Hi Sayan,

The OpenACC "cache" directive specifies which arrays should be put into shared memory. The cache directive will be available in the 12.7 release.

Here's an example:
Code:
void
testgpu(float *restrict *restrict b,float *restrict *restrict a,int *restrict ix,int *restrict iy,float *restrict t, int c)
{
    int i,j;
#pragma acc kernels copyin( a[0:N][0:M]), copyout(b[1:N-2][1:M-2])
    {
#pragma acc cache(a)
#pragma acc loop
        for( j = 1; j < M-1; ++j ){
            for( i = 1; i < N-1; ++i ){
                b[i][j] = 0.5f*(a[ix[i-1]][iy[j-1]] + a[i][j])/c;
            }
        }
    }

}


- Mat
Back to top
View user's profile
_sayan_



Joined: 07 Apr 2012
Posts: 29

PostPosted: Tue Jun 12, 2012 2:21 pm    Post subject: Reply with quote

Thank you Mat, does the compiler internally copy data into shared memory on certain cases, or acc directives at present only load data to/from registers or global memory?
Back to top
View user's profile
mkcolg



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

PostPosted: Tue Jun 12, 2012 4:33 pm    Post subject: Reply with quote

Hi Sayan,

Sorry, but I'm not entirely sure I understand the question. Do you mean does the compiler automatically find opportunities for shared memory usage or will shared memory only be utilized if the cache clause is explicitly set by the user?

Right now the PGI Accelerator Model with automatically use shared memory however OpenACC will require the user to explicitly set the cache directive. This is a difference between the two models.

- Mat
Back to top
View user's profile
_sayan_



Joined: 07 Apr 2012
Posts: 29

PostPosted: Wed Jun 13, 2012 8:28 am    Post subject: Reply with quote

Yes, my question was whether the compiler will find opportunities to use GPU shared memory. Thank you again for distinguishing the two models.
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