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  }