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

/* Test behavior of 'firstprivate' lexically vs. dynamically nested inside a
   'data' region.  */

#include <stdlib.h>


#define VERIFY(x) \
  do { \
    if (!(x)) \
      abort (); \
  } while (0);


/* This is basically and extended version of 't2' from 'firstprivate-1.c'.  */

int lexically_nested_val = 2;

static void
lexically_nested ()
{
#pragma acc data \
  copy (lexically_nested_val)
  {
    VERIFY (lexically_nested_val == 2);

#pragma acc parallel \
  present (lexically_nested_val)
    {
      VERIFY (lexically_nested_val == 2);

      /* This updates the device copy, or shared variable.  */
      lexically_nested_val = 7;
    }

#if ACC_MEM_SHARED
    VERIFY (lexically_nested_val == 7);
#else
    VERIFY (lexically_nested_val == 2);
#endif

    /* This only updates the local/shared variable, but not the device
       copy.  */
    lexically_nested_val = 5;

#pragma acc parallel \
    firstprivate (lexically_nested_val)
      {
#if 1 /* Current behavior.  */
	/* The 'firstprivate' copy is initialized from the device copy, or
	   shared variable.  */
# if ACC_MEM_SHARED
	VERIFY (lexically_nested_val == 5);
# else
	VERIFY (lexically_nested_val == 7);
# endif
#else /* Expected behavior per PR92036.  */
	/* The 'firstprivate' copy is initialized from the local thread.  */
	VERIFY (lexically_nested_val == 5);
#endif

	/* This updates the 'firstprivate' copy only, but not the shared
	   variable.  */
	lexically_nested_val = 9;
      }

    VERIFY (lexically_nested_val == 5);
  }
  /* If not shared, the device copy has now been copied back.  */

#if ACC_MEM_SHARED
  VERIFY (lexically_nested_val == 5);
#else
  VERIFY (lexically_nested_val == 7);
#endif
}


int dynamically_nested_val = 2;

/* Same as above, but compute construct 1 broken out, so no longer lexically
   nested inside 'data' region.  */

static void
dynamically_nested_compute_1 ()
{
#pragma acc parallel \
  present (dynamically_nested_val)
  {
    VERIFY (dynamically_nested_val == 2);

    /* This updates the device copy, or shared variable.  */
    dynamically_nested_val = 7;
  }
}

/* Same as above, but compute construct 2 broken out, so no longer lexically
   nested inside 'data' region.  */

static void
dynamically_nested_compute_2 ()
{
#pragma acc parallel \
  firstprivate (dynamically_nested_val)
    {
#if 1 /* Current behavior.  */
      /* The 'firstprivate' copy is initialized from the device copy, or shared
	 variable.  */
# if ACC_MEM_SHARED
      VERIFY (dynamically_nested_val == 5);
# else
      VERIFY (dynamically_nested_val == 7);
# endif
#else /* Expected behavior per PR92036.  */
      /* The 'firstprivate' copy is initialized from the local thread.  */
      VERIFY (dynamically_nested_val == 5);
#endif

      /* This updates the 'firstprivate' copy only, but not the shared
	 variable.  */
      dynamically_nested_val = 9;
    }
}

static void
dynamically_nested ()
{
#pragma acc data \
  copy (dynamically_nested_val)
  {
    VERIFY (dynamically_nested_val == 2);

    dynamically_nested_compute_1 ();

#if ACC_MEM_SHARED
    VERIFY (dynamically_nested_val == 7);
#else
    VERIFY (dynamically_nested_val == 2);
#endif

    /* This only updates the local/shared variable, but not the device
       copy.  */
    dynamically_nested_val = 5;

    dynamically_nested_compute_2 ();

    VERIFY (dynamically_nested_val == 5);
  }
  /* If not shared, the device copy has now been copied back.  */

#if ACC_MEM_SHARED
  VERIFY (dynamically_nested_val == 5);
#else
  VERIFY (dynamically_nested_val == 7);
#endif
}


int
main()
{
  lexically_nested ();
  dynamically_nested ();

  return 0;
}