Hi,

This a re-posting of the original note incorporating the suggestions
from Joseph and Nathan (thank you).

     This patch adds the processing of OpenACC host_data construct in C
     and C++. (Note: Support in Fortran is already in trunk.) The patch
     also adds the required support in the middle-end and libgomp.

     Background
         The host data construct is used to make an address of device
         data available on the host.

         The following illustrates use of the host data construct in
         conjunction with arrays which are already device-resident
         and an accelerator-only function.

                 int main(int argc, char **argv)
                 {
                   float *x, *y;
                   const int n = 1024;
                   int i;

                   x = (float*) malloc (n * sizeof(float));
                   y = (float*) malloc (n * sizeof(float));

                   /* Copy the arrays out to the device. */
                   #pragma acc data create(x[0:n]) copyout(y[0:n])
                   {
                     #pragma acc parallel
                     {
                       for (i = 0; i < n; i++)
                         {
                           x[i] = 1.0f;
                           y[i] = 0.0f;
                         }
                     }

                     /*
                      * The arrays are already on the device, so
                      * pass the device addresses to saxpy. NOTE:
                      * saxpy has been previously defined as an
                      * accelerator function.
                      */
                     #pragma acc host_data use_device(x, y)
                     {
                       saxpy(n, 2.0, x, 1, y, 1);
                     }
                   }

                   fprintf(stdout, "y[0] = %f\n", y[0]);
                   return 0;
                 }


     C and C++ front-ends

         Definitions for use by C and C++ were added to identify the
         host_data construct pragma and its' only valid clause: use_device.

         New functionality was added to do the parsing of the host_data
         pragma and validate the sole clause valid clause: use_device.
         As the host_data construct has associated with it a structured
         block, new functionality was added to build the compound
         statement to represent the block.

     Middle-end

         A gimple definition: GOVD_USE_DEVICE, has been added to indicate
         the use of the use_device clause. This flag is asserted as part
         of installing mappings into a omp context. The flag is subsequently
         reacted to during the gimplying of the host_data region's body.
         When this flag is encountered, an GOACC_deviceptr builtin call
         is inserted at the appropriate place.

     libgomp

         A new function has been added to handle pointer lookup for host
         data regions. As the comment in the code describes, this function
         will return the appropriate address based on whether it is called
         for the host or the target. This function is used in response to
         usage of the use_device clause.

     Tests

         New compile and runtime tests have been added.

     All of the code is in the gomp-4_0-branch.

     Regtested on x86_64-linux.

     Thanks!
     Jim