PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Basic Accelerator trouble with dynamically alloc arrays [C]

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



Joined: 07 Oct 2010
Posts: 1

PostPosted: Thu Oct 07, 2010 7:19 am    Post subject: Basic Accelerator trouble with dynamically alloc arrays [C] Reply with quote

Hi all,

In experimenting with the accelerator directives, we were unable to get a simple, "hello world" style program to compile (with the properly acclerated loop). The problem seems to lie with our dynamically allocated arrays, although I assumed this would be handled by the copyin() and copyout() mechanism. Here is the code:

Code:

#include <stdlib.h>                                                             
                                                                               
// compile with: pgcc test01.c -ta=nvidia -Minfo                               
                                                                               
int main (int argc, char** argv)                                               
{                                                                               
    int i;                                                                     
    float *a, *b;                                                               
    float c[1000];                                                             
                                                                               
    a = (float*) malloc (1000 * sizeof (float));                               
    b = (float*) malloc (1000 * sizeof (float));                               
                                                                               
    for (i = 0; i < 1000; i++)                                                 
        b[i] = 1;                                                             
   
    //this region does not parallize!?                                                                           
    #pragma acc region copyin(b[0:999]) copyout(a[0:999])                       
    {                                                                           
        for (i = 0; i < 1000; i++)                                             
            a[i] = b[i] + 1;                                                   
    }                                                                           
    //this region does parallize??                                                                           
    #pragma acc region                                                         
    {                                                                           
        for (i = 0; i < 1000; i++)                                             
            c[i] = a[i] + 1;                                                   
    }                                                                           
                                                                               
    return EXIT_SUCCESS;                                                       
}


When compiling with:
Code:

pgcc test01.c -ta=nvidia -Minfo

We get the following error:
Code:

main:                                                                           
     14, Memory set idiom, loop replaced by call to __c_mset4                   
     17, No parallel kernels found, accelerator region ignored                 
     19, Complex loop carried dependence of 'b' prevents parallelization       
         Loop carried dependence of 'a' prevents parallelization               
         Loop carried backward dependence of 'a' prevents vectorization         
     23, Generating copyin(a[0:999])                                           
         Generating copyout(c[0:999])                                           
         Generating compute capability 1.0 binary                               
         Generating compute capability 1.3 binary                               
     25, Loop is parallelizable                                                 
         Accelerator kernel generated                                           
         25, #pragma acc for parallel, vector(256)                             
             CC 1.0 : 3 registers; 20 shared, 28 constant, 0 local memory bytes\
; 100 occupancy                                                                 
             CC 1.3 : 3 registers; 20 shared, 28 constant, 0 local memory bytes\
; 100 occupancy


Am I crazy or just blind? It seems like it should work assuming you give it the proper copy() directive to tell the compiler how big the variables are.
Back to top
View user's profile
mkcolg



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

PostPosted: Sat Oct 09, 2010 1:22 am    Post subject: Reply with quote

Hi maximilian,

You need to declare your pointers using the C99 restrict keyword or compile with "-Msafeptr". Otherwise, the compiler must assume that pointers overlap and hence cannot be parallelzied.

Code:
 float * restrict a, * restrict b;


Also, you may wish to use the flag "-Mfcon" or append "f" to your constant real values to make them single precision. By default, constant reals are double precision.

Hope this helps,
Mat
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