libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c

   1 /* Test behavior of 'firstprivate' lexically vs. dynamically nested inside a
   2    'data' region.  */
   3
   4 #include <stdlib.h>
   5
   6
   7 #define VERIFY(x) \
   8   do { \
   9     if (!(x)) \
  10       abort (); \
  11   } while (0);
  12
  13
  14 /* This is basically and extended version of 't2' from 'firstprivate-1.c'.  */
  15
  16 int lexically_nested_val = 2;
  17
  18 static void
  19 lexically_nested ()
  20 {
  21 #pragma acc data \
  22   copy (lexically_nested_val)
  23   {
  24     VERIFY (lexically_nested_val == 2);
  25
  26 #pragma acc parallel \
  27   present (lexically_nested_val)
  28     {
  29       VERIFY (lexically_nested_val == 2);
  30
  31       /* This updates the device copy, or shared variable.  */
  32       lexically_nested_val = 7;
  33     }
  34
  35 #if ACC_MEM_SHARED
  36     VERIFY (lexically_nested_val == 7);
  37 #else
  38     VERIFY (lexically_nested_val == 2);
  39 #endif
  40
  41     /* This only updates the local/shared variable, but not the device
  42        copy.  */
  43     lexically_nested_val = 5;
  44
  45 #pragma acc parallel \
  46     firstprivate (lexically_nested_val)
  47       {
  48 #if 1 /* Current behavior.  */
  49         /* The 'firstprivate' copy is initialized from the device copy, or
  50            shared variable.  */
  51 # if ACC_MEM_SHARED
  52         VERIFY (lexically_nested_val == 5);
  53 # else
  54         VERIFY (lexically_nested_val == 7);
  55 # endif
  56 #else /* Expected behavior per PR92036.  */
  57         /* The 'firstprivate' copy is initialized from the local thread.  */
  58         VERIFY (lexically_nested_val == 5);
  59 #endif
  60
  61         /* This updates the 'firstprivate' copy only, but not the shared
  62            variable.  */
  63         lexically_nested_val = 9;
  64       }
  65
  66     VERIFY (lexically_nested_val == 5);
  67   }
  68   /* If not shared, the device copy has now been copied back.  */
  69
  70 #if ACC_MEM_SHARED
  71   VERIFY (lexically_nested_val == 5);
  72 #else
  73   VERIFY (lexically_nested_val == 7);
  74 #endif
  75 }
  76
  77
  78 int dynamically_nested_val = 2;
  79
  80 /* Same as above, but compute construct 1 broken out, so no longer lexically
  81    nested inside 'data' region.  */
  82
  83 static void
  84 dynamically_nested_compute_1 ()
  85 {
  86 #pragma acc parallel \
  87   present (dynamically_nested_val)
  88   {
  89     VERIFY (dynamically_nested_val == 2);
  90
  91     /* This updates the device copy, or shared variable.  */
  92     dynamically_nested_val = 7;
  93   }
  94 }
  95
  96 /* Same as above, but compute construct 2 broken out, so no longer lexically
  97    nested inside 'data' region.  */
  98
  99 static void
 100 dynamically_nested_compute_2 ()
 101 {
 102 #pragma acc parallel \
 103   firstprivate (dynamically_nested_val)
 104     {
 105 #if 1 /* Current behavior.  */
 106       /* The 'firstprivate' copy is initialized from the device copy, or shared
 107          variable.  */
 108 # if ACC_MEM_SHARED
 109       VERIFY (dynamically_nested_val == 5);
 110 # else
 111       VERIFY (dynamically_nested_val == 7);
 112 # endif
 113 #else /* Expected behavior per PR92036.  */
 114       /* The 'firstprivate' copy is initialized from the local thread.  */
 115       VERIFY (dynamically_nested_val == 5);
 116 #endif
 117
 118       /* This updates the 'firstprivate' copy only, but not the shared
 119          variable.  */
 120       dynamically_nested_val = 9;
 121     }
 122 }
 123
 124 static void
 125 dynamically_nested ()
 126 {
 127 #pragma acc data \
 128   copy (dynamically_nested_val)
 129   {
 130     VERIFY (dynamically_nested_val == 2);
 131
 132     dynamically_nested_compute_1 ();
 133
 134 #if ACC_MEM_SHARED
 135     VERIFY (dynamically_nested_val == 7);
 136 #else
 137     VERIFY (dynamically_nested_val == 2);
 138 #endif
 139
 140     /* This only updates the local/shared variable, but not the device
 141        copy.  */
 142     dynamically_nested_val = 5;
 143
 144     dynamically_nested_compute_2 ();
 145
 146     VERIFY (dynamically_nested_val == 5);
 147   }
 148   /* If not shared, the device copy has now been copied back.  */
 149
 150 #if ACC_MEM_SHARED
 151   VERIFY (dynamically_nested_val == 5);
 152 #else
 153   VERIFY (dynamically_nested_val == 7);
 154 #endif
 155 }
 156
 157
 158 int
 159 main()
 160 {
 161   lexically_nested ();
 162   dynamically_nested ();
 163
 164   return 0;
 165 }