From: Thomas Schwinge Date: Tue, 18 Jun 2019 22:15:03 +0000 (+0200) Subject: Test cases to verify OpenACC 'firstprivate' mappings X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=85fca03a09ababbe8808cbe402368e375b8834b9;p=gcc.git Test cases to verify OpenACC 'firstprivate' mappings gcc/testsuite/ * c-c++-common/goacc/firstprivate-mappings-1.c: New file. * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: New file. * testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c: Likewise. From-SVN: r272451 --- diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 699a94b3ed4..bffa899c767 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,5 +1,8 @@ 2019-06-18 Thomas Schwinge + * c-c++-common/goacc/firstprivate-mappings-1.c: New file. + * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. + PR testsuite/90861 * c-c++-common/goacc/declare-pr90861.c: New file. diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c new file mode 100644 index 00000000000..c8270472a9c --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c @@ -0,0 +1,533 @@ +/* Verify OpenACC 'firstprivate' mappings. */ + +/* This file is also sourced from + '../../../../libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c' + as an execution test. */ + +/* See also '../../g++.dg/goacc/firstprivate-mappings-1.C'. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +/* { dg-additional-options "-fext-numeric-literals" { target c++ } } */ + +/* { dg-additional-options "-Wno-psabi" } as apparently we're doing funny + things with vector arguments. */ + +#include +#include +#include + + +#ifdef __SIZEOF_INT128__ +# define HAVE_INT128 1 +#else +# define HAVE_INT128 0 +#endif + + +/* The one is only relevant for offloading compilation; will always be enabled + when doing tree scanning. */ +#ifdef ACC_DEVICE_TYPE_nvidia +/* PR71064. */ +# define DO_LONG_DOUBLE 0 +#else +# define DO_LONG_DOUBLE 1 +#endif + + +/* Simplify scanning for function names in tree dumps. */ +#ifdef __cplusplus +extern "C" { +#endif + + +/* Inside the following OpenACC 'parallel' constructs' regions, we modify the + 'firstprivate' variables, so that we can check that we don't copy these + back. */ + + +static void +p (short *spi) +{ + short *spo; +#pragma acc parallel \ + copyout (spo) \ + firstprivate (spi) + { + spo = ++spi; + } + if (spo != spi + 1) + __builtin_abort (); +} + + +static void +b (bool bi) +{ + bool bo; +#pragma acc parallel \ + copyout (bo) \ + firstprivate (bi) + { + bo = (bi = !bi); + } + if (bo != !bi) + __builtin_abort (); +} + + +static void +i (int8_t i8i, + uint8_t u8i, + int16_t i16i, + uint16_t u16i, + int32_t i32i, + uint32_t u32i, + int64_t i64i, + uint64_t u64i) +{ + int8_t i8o; + uint8_t u8o; + int16_t i16o; + uint16_t u16o; + int32_t i32o; + uint32_t u32o; + int64_t i64o; + uint64_t u64o; +#pragma acc parallel \ + copyout (i8o) \ + firstprivate (i8i) \ + copyout (u8o) \ + firstprivate (u8i) \ + copyout (i16o) \ + firstprivate (i16i) \ + copyout (u16o) \ + firstprivate (u16i) \ + copyout (i32o) \ + firstprivate (i32i) \ + copyout (u32o) \ + firstprivate (u32i) \ + copyout (i64o) \ + firstprivate (i64i) \ + copyout (u64o) \ + firstprivate (u64i) + { + i8o = --i8i; + u8o = ++u8i; + i16o = --i16i; + u16o = ++u16i; + i32o = --i32i; + u32o = ++u32i; + i64o = --i64i; + u64o = ++u64i; + } + if (i8o != i8i - 1) + __builtin_abort (); + if (u8o != u8i + 1) + __builtin_abort (); + if (i16o != i16i - 1) + __builtin_abort (); + if (u16o != u16i + 1) + __builtin_abort (); + if (i32o != i32i - 1) + __builtin_abort (); + if (u32o != u32i + 1) + __builtin_abort (); + if (i64o != i64i - 1) + __builtin_abort (); + if (u64o != u64i + 1) + __builtin_abort (); +} + + +#if HAVE_INT128 +static void +i128 (__int128 i128i, unsigned __int128 u128i) +{ + __int128 i128o; + unsigned __int128 u128o; +# pragma acc parallel \ + copyout (i128o) \ + firstprivate (i128i) \ + copyout(u128o) \ + firstprivate (u128i) + { + i128o = --i128i; + u128o = ++u128i; + } + if (i128o != i128i - 1) + __builtin_abort (); + if (u128o != u128i + 1) + __builtin_abort (); +} +#endif + + +static void +flt_dbl (float flti, double dbli) +{ + float flto; + double dblo; +#pragma acc parallel \ + copyout (flto) \ + firstprivate (flti) \ + copyout (dblo) \ + firstprivate (dbli) + { + flto = --flti; + dblo = --dbli; + } + if (flto != flti - 1) + __builtin_abort (); + if (dblo != dbli - 1) + __builtin_abort (); +} + + +static void +ldbl (long double ldbli) +{ +#if DO_LONG_DOUBLE + long double ldblo; +# pragma acc parallel \ + copyout (ldblo) \ + firstprivate (ldbli) + { + ldblo = --ldbli; + } + if (ldblo != ldbli - 1) + __builtin_abort (); +#endif +} + + +static void +c (_Complex unsigned char cuci, + _Complex signed short cssi, + _Complex unsigned int cuii, + _Complex signed long csli, + _Complex float cflti, + _Complex double cdbli) +{ + _Complex unsigned char cuco; + _Complex signed short csso; + _Complex unsigned int cuio; + _Complex signed long cslo; + _Complex float cflto; + _Complex double cdblo; +#pragma acc parallel \ + copyout (cuco) \ + firstprivate (cuci) \ + copyout (csso) \ + firstprivate (cssi) \ + copyout (cuio) \ + firstprivate (cuii) \ + copyout (cslo) \ + firstprivate (csli) \ + copyout (cflto) \ + firstprivate (cflti) \ + copyout (cdblo) \ + firstprivate (cdbli) + { + cuco = (cuci += (1 + 1j)); + csso = (cssi -= (1 + 1j)); + cuio = (cuii += (1 + 1j)); + cslo = (csli -= (1 + 1j)); + cflto = (cflti -= (1 + 1j)); + cdblo = (cdbli -= (1 + 1j)); + } + if (cuco != cuci + (1 + 1j)) + __builtin_abort (); + if (csso != cssi - (1 + 1j)) + __builtin_abort (); + if (cuio != cuii + (1 + 1j)) + __builtin_abort (); + if (cslo != csli - (1 + 1j)) + __builtin_abort (); + if (cflto != cflti - (1 + 1j)) + __builtin_abort (); + if (cdblo != cdbli - (1 + 1j)) + __builtin_abort (); +} + + +static void +cldbl (_Complex long double cldbli) +{ +#if DO_LONG_DOUBLE + _Complex long double cldblo; +# pragma acc parallel \ + copyout (cldblo) \ + firstprivate (cldbli) + { + cldblo = (cldbli -= (1 + 1j)); + } + if (cldblo != cldbli - (1 + 1j)) + __builtin_abort (); +#endif +} + + +#define V_EQ(v1, v2) \ + ({ \ + __typeof__ (v1) v_d = (v1) != (v2); \ + __typeof__ (v_d) v_0 = { 0 }; \ + memcmp (&v_d, &v_0, sizeof v_d) == 0; \ + }) + +typedef uint8_t __attribute__ ((vector_size (2 * sizeof (uint8_t)))) v2u8; +typedef int16_t __attribute__ ((vector_size (4 * sizeof (int16_t)))) v4i16; +typedef uint32_t __attribute__ ((vector_size (8 * sizeof (uint32_t)))) v8u32; +typedef int64_t __attribute__ ((vector_size (16 * sizeof (int64_t)))) v16i64; +typedef float __attribute__ ((vector_size (1 * sizeof (float)))) v1flt; +typedef float __attribute__ ((vector_size (2 * sizeof (float)))) v2flt; +typedef float __attribute__ ((vector_size (4 * sizeof (float)))) v4flt; +typedef float __attribute__ ((vector_size (8 * sizeof (float)))) v8flt; +typedef double __attribute__ ((vector_size (1 * sizeof (double)))) v1dbl; +typedef double __attribute__ ((vector_size (2 * sizeof (double)))) v2dbl; +typedef double __attribute__ ((vector_size (4 * sizeof (double)))) v4dbl; +typedef double __attribute__ ((vector_size (8 * sizeof (double)))) v8dbl; + +static void +v (v2u8 v2u8i, v4i16 v4i16i, v8u32 v8u32i, v16i64 v16i64i, + v1flt v1flti, v2flt v2flti, v4flt v4flti, v8flt v8flti, + v1dbl v1dbli, v2dbl v2dbli, v4dbl v4dbli, v8dbl v8dbli) +{ + v2u8 v2u8o; + v4i16 v4i16o; + v8u32 v8u32o; + v16i64 v16i64o; + v1flt v1flto; + v2flt v2flto; + v4flt v4flto; + v8flt v8flto; + v1dbl v1dblo; + v2dbl v2dblo; + v4dbl v4dblo; + v8dbl v8dblo; +#pragma acc parallel \ + copyout (v2u8o) \ + firstprivate (v2u8i) \ + copyout (v4i16o) \ + firstprivate (v4i16i) \ + copyout (v8u32o) \ + firstprivate (v8u32i) \ + copyout (v16i64o) \ + firstprivate (v16i64i) \ + copyout (v1flto) \ + firstprivate (v1flti) \ + copyout (v2flto) \ + firstprivate (v2flti) \ + copyout (v4flto) \ + firstprivate (v4flti) \ + copyout (v8flto) \ + firstprivate (v8flti) \ + copyout (v1dblo) \ + firstprivate (v1dbli) \ + copyout (v2dblo) \ + firstprivate (v2dbli) \ + copyout (v4dblo) \ + firstprivate (v4dbli) \ + copyout (v8dblo) \ + firstprivate (v8dbli) + { + v2u8o = ++v2u8i; + v4i16o = --v4i16i; + v8u32o = ++v8u32i; + v16i64o = --v16i64i; + v1flto = --v1flti; + v2flto = --v2flti; + v4flto = --v4flti; + v8flto = --v8flti; + v1dblo = --v1dbli; + v2dblo = --v2dbli; + v4dblo = --v4dbli; + v8dblo = --v8dbli; + } + if (!V_EQ (v2u8o, v2u8i + 1)) + __builtin_abort (); + if (!V_EQ (v4i16o, v4i16i - 1)) + __builtin_abort (); + if (!V_EQ (v8u32o, v8u32i + 1)) + __builtin_abort (); + if (!V_EQ (v16i64o, v16i64i - 1)) + __builtin_abort (); + if (!V_EQ (v1flto, v1flti - 1)) + __builtin_abort (); + if (!V_EQ (v2flto, v2flti - 1)) + __builtin_abort (); + if (!V_EQ (v4flto, v4flti - 1)) + __builtin_abort (); + if (!V_EQ (v8flto, v8flti - 1)) + __builtin_abort (); + if (!V_EQ (v1dblo, v1dbli - 1)) + __builtin_abort (); + if (!V_EQ (v2dblo, v2dbli - 1)) + __builtin_abort (); + if (!V_EQ (v4dblo, v4dbli - 1)) + __builtin_abort (); + if (!V_EQ (v8dblo, v8dbli - 1)) + __builtin_abort (); +} + + +/* "error: could not find an integer type of the same size as 'long double'" */ +#if HAVE_INT128 +typedef long double __attribute__ ((vector_size (1 * sizeof (long double)))) v1ldbl; +typedef long double __attribute__ ((vector_size (2 * sizeof (long double)))) v2ldbl; +typedef long double __attribute__ ((vector_size (4 * sizeof (long double)))) v4ldbl; +typedef long double __attribute__ ((vector_size (8 * sizeof (long double)))) v8ldbl; + +static void +vldbl (v1ldbl v1ldbli, v2ldbl v2ldbli, v4ldbl v4ldbli, v8ldbl v8ldbli) +{ +# if DO_LONG_DOUBLE + v1ldbl v1ldblo; + v2ldbl v2ldblo; + v4ldbl v4ldblo; + v8ldbl v8ldblo; +# pragma acc parallel \ + copyout (v1ldblo) \ + firstprivate (v1ldbli) \ + copyout (v2ldblo) \ + firstprivate (v2ldbli) \ + copyout (v4ldblo) \ + firstprivate (v4ldbli) \ + copyout (v8ldblo) \ + firstprivate (v8ldbli) + { + v1ldblo = --v1ldbli; + v2ldblo = --v2ldbli; + v4ldblo = --v4ldbli; + v8ldblo = --v8ldbli; + } + if (!V_EQ (v1ldblo, v1ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v2ldblo, v2ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v4ldblo, v4ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v8ldblo, v8ldbli - 1)) + __builtin_abort (); +# endif +} +#endif + + +static void +vla (int array_li) +{ + _Complex double array[array_li]; + uint32_t array_so; +#pragma acc parallel \ + copyout (array_so) + /* The gimplifier has created an implicit 'firstprivate' clause for the array + length. + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } } + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } } } + (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */ + { + array_so = sizeof array; + } + if (array_so != sizeof array) + __builtin_abort (); +} + + +#ifdef __cplusplus +} +#endif + + +int +main (int argc, char *argv[]) +{ + { + short s; + short *sp = &s; + p (sp); + } + + { + bool bi = true; + b (bi); + } + + { + int8_t i8i = -1; + uint8_t u8i = 1; + int16_t i16i = -2; + uint16_t u16i = 2; + int32_t i32i = -3; + uint32_t u32i = 3; + int64_t i64i = -4; + uint64_t u64i = 4; + i (i8i, u8i, i16i, u16i, i32i, u32i, i64i, u64i); + } + +#if HAVE_INT128 + { + __int128 i128i = -8; + unsigned __int128 u128i = 8; + i128 (i128i, u128i); + } +#endif + + { + float flti = .5; + double dbli = .25; + flt_dbl (flti, dbli); + } + + { + long double ldbli = .125; + ldbl (ldbli); + } + + { + _Complex unsigned char cuci = 1 + 2j; + _Complex signed short cssi = -2 + (-4j); + _Complex unsigned int cuii = 3 + 6j; + _Complex signed long csli = -4 + (-8j); + _Complex float cflti = .5 + 1j; + _Complex double cdbli = .25 + .5j; + c (cuci, cssi, cuii, csli, cflti, cdbli); + } + + { + _Complex long double cldbli = .125 + .25j; + cldbl (cldbli); + } + + { + v2u8 v2u8i = {2, 3}; + v4i16 v4i16i = { -1, -2, 5, 4 }; + v8u32 v8u32i = { 3, 6, 9, 11}; + v16i64 v16i64i = { 10, 21, -25, 44, 31, -1, 1, 222, -1, -12, 52, -44, -13, 1, -1, -222}; + v1flt v1flti = { -.5 }; + v2flt v2flti = { 1.5, -2.5 }; + v4flt v4flti = { 3.5, -4.5, -5.5, -6.5 }; + v8flt v8flti = { -7.5, 8.5, 9.5, 10.5, -11.5, -12.5, 13.5, 14.5 }; + v1dbl v1dbli = { 0.25 }; + v2dbl v2dbli = { -1.25, -2.25 }; + v4dbl v4dbli = { 3.25, -4.25, 5.25, 6.25 }; + v8dbl v8dbli = { 7.25, 8.25, -9.25, -10.25, -11.25, 12.25, 13.25, -14.25 }; + v (v2u8i, v4i16i, v8u32i, v16i64i, + v1flti, v2flti, v4flti, v8flti, + v1dbli, v2dbli, v4dbli, v8dbli); + } + +#if HAVE_INT128 + { + v1ldbl v1ldbli = { -0.125 }; + v2ldbl v2ldbli = { 1.125, -2.125 }; + v4ldbl v4ldbli = { -3.125, -4.125, 5.125, -6.125 }; + v8ldbl v8ldbli = { 7.125, -8.125, -9.125, 10.125, 11.125, 12.125, 13.125, 14.125 }; + vldbl (v1ldbli, v2ldbli, v4ldbli, v8ldbli); + } +#endif + + vla (argc); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C new file mode 100644 index 00000000000..639bf3f3299 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C @@ -0,0 +1,529 @@ +/* Verify OpenACC 'firstprivate' mappings for C++ reference types. */ + +/* This file is also sourced from + '../../../../libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C' + as an execution test. */ + +/* See also '../../c-c++-common/goacc/firstprivate-mappings-1.c'. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +/* { dg-additional-options "-fext-numeric-literals" { target c++ } } */ + +#include +#include +#include + + +#ifdef __SIZEOF_INT128__ +# define HAVE_INT128 1 +#else +# define HAVE_INT128 0 +#endif + + +/* The one is only relevant for offloading compilation; will always be enabled + when doing tree scanning. */ +#ifdef ACC_DEVICE_TYPE_nvidia +/* PR71064. */ +# define DO_LONG_DOUBLE 0 +#else +# define DO_LONG_DOUBLE 1 +#endif + + +/* Simplify scanning for function names in tree dumps. */ +#ifdef __cplusplus +extern "C" { +#endif + + +/* Inside the following OpenACC 'parallel' constructs' regions, we modify the + 'firstprivate' variables, so that we can check that we don't copy these + back. */ + + +static void +p (short *&spi) +{ + short *spo; +#pragma acc parallel \ + copyout (spo) \ + firstprivate (spi) + { + spo = ++spi; + } + if (spo != spi + 1) + __builtin_abort (); +} + + +static void +b (bool &bi) +{ + bool bo; +#pragma acc parallel \ + copyout (bo) \ + firstprivate (bi) + { + bo = (bi = !bi); + } + if (bo != !bi) + __builtin_abort (); +} + + +static void +i (int8_t &i8i, + uint8_t &u8i, + int16_t &i16i, + uint16_t &u16i, + int32_t &i32i, + uint32_t &u32i, + int64_t &i64i, + uint64_t &u64i) +{ + int8_t i8o; + uint8_t u8o; + int16_t i16o; + uint16_t u16o; + int32_t i32o; + uint32_t u32o; + int64_t i64o; + uint64_t u64o; +#pragma acc parallel \ + copyout (i8o) \ + firstprivate (i8i) \ + copyout (u8o) \ + firstprivate (u8i) \ + copyout (i16o) \ + firstprivate (i16i) \ + copyout (u16o) \ + firstprivate (u16i) \ + copyout (i32o) \ + firstprivate (i32i) \ + copyout (u32o) \ + firstprivate (u32i) \ + copyout (i64o) \ + firstprivate (i64i) \ + copyout (u64o) \ + firstprivate (u64i) + { + i8o = --i8i; + u8o = ++u8i; + i16o = --i16i; + u16o = ++u16i; + i32o = --i32i; + u32o = ++u32i; + i64o = --i64i; + u64o = ++u64i; + } + if (i8o != i8i - 1) + __builtin_abort (); + if (u8o != u8i + 1) + __builtin_abort (); + if (i16o != i16i - 1) + __builtin_abort (); + if (u16o != u16i + 1) + __builtin_abort (); + if (i32o != i32i - 1) + __builtin_abort (); + if (u32o != u32i + 1) + __builtin_abort (); + if (i64o != i64i - 1) + __builtin_abort (); + if (u64o != u64i + 1) + __builtin_abort (); +} + + +#if HAVE_INT128 +static void +i128 (__int128 &i128i, unsigned __int128 &u128i) +{ + __int128 i128o; + unsigned __int128 u128o; +# pragma acc parallel \ + copyout (i128o) \ + firstprivate (i128i) \ + copyout(u128o) \ + firstprivate (u128i) + { + i128o = --i128i; + u128o = ++u128i; + } + if (i128o != i128i - 1) + __builtin_abort (); + if (u128o != u128i + 1) + __builtin_abort (); +} +#endif + + +static void +flt_dbl (float &flti, double &dbli) +{ + float flto; + double dblo; +#pragma acc parallel \ + copyout (flto) \ + firstprivate (flti) \ + copyout (dblo) \ + firstprivate (dbli) + { + flto = --flti; + dblo = --dbli; + } + if (flto != flti - 1) + __builtin_abort (); + if (dblo != dbli - 1) + __builtin_abort (); +} + + +static void +ldbl (long double &ldbli) +{ +#if DO_LONG_DOUBLE + long double ldblo; +# pragma acc parallel \ + copyout (ldblo) \ + firstprivate (ldbli) + { + ldblo = --ldbli; + } + if (ldblo != ldbli - 1) + __builtin_abort (); +#endif +} + + +static void +c (_Complex unsigned char &cuci, + _Complex signed short &cssi, + _Complex unsigned int &cuii, + _Complex signed long &csli, + _Complex float &cflti, + _Complex double &cdbli) +{ + _Complex unsigned char cuco; + _Complex signed short csso; + _Complex unsigned int cuio; + _Complex signed long cslo; + _Complex float cflto; + _Complex double cdblo; +#pragma acc parallel \ + copyout (cuco) \ + firstprivate (cuci) \ + copyout (csso) \ + firstprivate (cssi) \ + copyout (cuio) \ + firstprivate (cuii) \ + copyout (cslo) \ + firstprivate (csli) \ + copyout (cflto) \ + firstprivate (cflti) \ + copyout (cdblo) \ + firstprivate (cdbli) + { + cuco = (cuci += (1 + 1j)); + csso = (cssi -= (1 + 1j)); + cuio = (cuii += (1 + 1j)); + cslo = (csli -= (1 + 1j)); + cflto = (cflti -= (1 + 1j)); + cdblo = (cdbli -= (1 + 1j)); + } + if (cuco != cuci + (1 + 1j)) + __builtin_abort (); + if (csso != cssi - (1 + 1j)) + __builtin_abort (); + if (cuio != cuii + (1 + 1j)) + __builtin_abort (); + if (cslo != csli - (1 + 1j)) + __builtin_abort (); + if (cflto != cflti - (1 + 1j)) + __builtin_abort (); + if (cdblo != cdbli - (1 + 1j)) + __builtin_abort (); +} + + +static void +cldbl (_Complex long double &cldbli) +{ +#if DO_LONG_DOUBLE + _Complex long double cldblo; +# pragma acc parallel \ + copyout (cldblo) \ + firstprivate (cldbli) + { + cldblo = (cldbli -= (1 + 1j)); + } + if (cldblo != cldbli - (1 + 1j)) + __builtin_abort (); +#endif +} + + +#define V_EQ(v1, v2) \ + ({ \ + __typeof__ (v1) v_d = (v1) != (v2); \ + __typeof__ (v_d) v_0 = { 0 }; \ + memcmp (&v_d, &v_0, sizeof v_d) == 0; \ + }) + +typedef uint8_t __attribute__ ((vector_size (2 * sizeof (uint8_t)))) v2u8; +typedef int16_t __attribute__ ((vector_size (4 * sizeof (int16_t)))) v4i16; +typedef uint32_t __attribute__ ((vector_size (8 * sizeof (uint32_t)))) v8u32; +typedef int64_t __attribute__ ((vector_size (16 * sizeof (int64_t)))) v16i64; +typedef float __attribute__ ((vector_size (1 * sizeof (float)))) v1flt; +typedef float __attribute__ ((vector_size (2 * sizeof (float)))) v2flt; +typedef float __attribute__ ((vector_size (4 * sizeof (float)))) v4flt; +typedef float __attribute__ ((vector_size (8 * sizeof (float)))) v8flt; +typedef double __attribute__ ((vector_size (1 * sizeof (double)))) v1dbl; +typedef double __attribute__ ((vector_size (2 * sizeof (double)))) v2dbl; +typedef double __attribute__ ((vector_size (4 * sizeof (double)))) v4dbl; +typedef double __attribute__ ((vector_size (8 * sizeof (double)))) v8dbl; + +static void +v (v2u8 &v2u8i, v4i16 &v4i16i, v8u32 &v8u32i, v16i64 &v16i64i, + v1flt &v1flti, v2flt &v2flti, v4flt &v4flti, v8flt &v8flti, + v1dbl &v1dbli, v2dbl &v2dbli, v4dbl &v4dbli, v8dbl &v8dbli) +{ + v2u8 v2u8o; + v4i16 v4i16o; + v8u32 v8u32o; + v16i64 v16i64o; + v1flt v1flto; + v2flt v2flto; + v4flt v4flto; + v8flt v8flto; + v1dbl v1dblo; + v2dbl v2dblo; + v4dbl v4dblo; + v8dbl v8dblo; +#pragma acc parallel \ + copyout (v2u8o) \ + firstprivate (v2u8i) \ + copyout (v4i16o) \ + firstprivate (v4i16i) \ + copyout (v8u32o) \ + firstprivate (v8u32i) \ + copyout (v16i64o) \ + firstprivate (v16i64i) \ + copyout (v1flto) \ + firstprivate (v1flti) \ + copyout (v2flto) \ + firstprivate (v2flti) \ + copyout (v4flto) \ + firstprivate (v4flti) \ + copyout (v8flto) \ + firstprivate (v8flti) \ + copyout (v1dblo) \ + firstprivate (v1dbli) \ + copyout (v2dblo) \ + firstprivate (v2dbli) \ + copyout (v4dblo) \ + firstprivate (v4dbli) \ + copyout (v8dblo) \ + firstprivate (v8dbli) + { + v2u8o = ++v2u8i; + v4i16o = --v4i16i; + v8u32o = ++v8u32i; + v16i64o = --v16i64i; + v1flto = --v1flti; + v2flto = --v2flti; + v4flto = --v4flti; + v8flto = --v8flti; + v1dblo = --v1dbli; + v2dblo = --v2dbli; + v4dblo = --v4dbli; + v8dblo = --v8dbli; + } + if (!V_EQ (v2u8o, v2u8i + 1)) + __builtin_abort (); + if (!V_EQ (v4i16o, v4i16i - 1)) + __builtin_abort (); + if (!V_EQ (v8u32o, v8u32i + 1)) + __builtin_abort (); + if (!V_EQ (v16i64o, v16i64i - 1)) + __builtin_abort (); + if (!V_EQ (v1flto, v1flti - 1)) + __builtin_abort (); + if (!V_EQ (v2flto, v2flti - 1)) + __builtin_abort (); + if (!V_EQ (v4flto, v4flti - 1)) + __builtin_abort (); + if (!V_EQ (v8flto, v8flti - 1)) + __builtin_abort (); + if (!V_EQ (v1dblo, v1dbli - 1)) + __builtin_abort (); + if (!V_EQ (v2dblo, v2dbli - 1)) + __builtin_abort (); + if (!V_EQ (v4dblo, v4dbli - 1)) + __builtin_abort (); + if (!V_EQ (v8dblo, v8dbli - 1)) + __builtin_abort (); +} + + +/* "error: could not find an integer type of the same size as 'long double'" */ +#if HAVE_INT128 +typedef long double __attribute__ ((vector_size (1 * sizeof (long double)))) v1ldbl; +typedef long double __attribute__ ((vector_size (2 * sizeof (long double)))) v2ldbl; +typedef long double __attribute__ ((vector_size (4 * sizeof (long double)))) v4ldbl; +typedef long double __attribute__ ((vector_size (8 * sizeof (long double)))) v8ldbl; + +static void +vldbl (v1ldbl &v1ldbli, v2ldbl &v2ldbli, v4ldbl &v4ldbli, v8ldbl &v8ldbli) +{ +# if DO_LONG_DOUBLE + v1ldbl v1ldblo; + v2ldbl v2ldblo; + v4ldbl v4ldblo; + v8ldbl v8ldblo; +# pragma acc parallel \ + copyout (v1ldblo) \ + firstprivate (v1ldbli) \ + copyout (v2ldblo) \ + firstprivate (v2ldbli) \ + copyout (v4ldblo) \ + firstprivate (v4ldbli) \ + copyout (v8ldblo) \ + firstprivate (v8ldbli) + { + v1ldblo = --v1ldbli; + v2ldblo = --v2ldbli; + v4ldblo = --v4ldbli; + v8ldblo = --v8ldbli; + } + if (!V_EQ (v1ldblo, v1ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v2ldblo, v2ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v4ldblo, v4ldbli - 1)) + __builtin_abort (); + if (!V_EQ (v8ldblo, v8ldbli - 1)) + __builtin_abort (); +# endif +} +#endif + + +static void +vla (int &array_li) +{ + _Complex double array[array_li]; + uint32_t array_so; +#pragma acc parallel \ + copyout (array_so) + /* The gimplifier has created an implicit 'firstprivate' clause for the array + length. + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } } + (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */ + { + array_so = sizeof array; + } + if (array_so != sizeof array) + __builtin_abort (); +} + + +#ifdef __cplusplus +} +#endif + + +int +main (int argc, char *argv[]) +{ + { + short s; + short *sp = &s; + p (sp); + } + + { + bool bi = true; + b (bi); + } + + { + int8_t i8i = -1; + uint8_t u8i = 1; + int16_t i16i = -2; + uint16_t u16i = 2; + int32_t i32i = -3; + uint32_t u32i = 3; + int64_t i64i = -4; + uint64_t u64i = 4; + i (i8i, u8i, i16i, u16i, i32i, u32i, i64i, u64i); + } + +#if HAVE_INT128 + { + __int128 i128i = -8; + unsigned __int128 u128i = 8; + i128 (i128i, u128i); + } +#endif + + { + float flti = .5; + double dbli = .25; + flt_dbl (flti, dbli); + } + + { + long double ldbli = .125; + ldbl (ldbli); + } + + { + _Complex unsigned char cuci = 1 + 2j; + _Complex signed short cssi = -2 + (-4j); + _Complex unsigned int cuii = 3 + 6j; + _Complex signed long csli = -4 + (-8j); + _Complex float cflti = .5 + 1j; + _Complex double cdbli = .25 + .5j; + c (cuci, cssi, cuii, csli, cflti, cdbli); + } + + { + _Complex long double cldbli = .125 + .25j; + cldbl (cldbli); + } + + { + v2u8 v2u8i = {2, 3}; + v4i16 v4i16i = { -1, -2, 5, 4 }; + v8u32 v8u32i = { 3, 6, 9, 11}; + v16i64 v16i64i = { 10, 21, -25, 44, 31, -1, 1, 222, -1, -12, 52, -44, -13, 1, -1, -222}; + v1flt v1flti = { -.5 }; + v2flt v2flti = { 1.5, -2.5 }; + v4flt v4flti = { 3.5, -4.5, -5.5, -6.5 }; + v8flt v8flti = { -7.5, 8.5, 9.5, 10.5, -11.5, -12.5, 13.5, 14.5 }; + v1dbl v1dbli = { 0.25 }; + v2dbl v2dbli = { -1.25, -2.25 }; + v4dbl v4dbli = { 3.25, -4.25, 5.25, 6.25 }; + v8dbl v8dbli = { 7.25, 8.25, -9.25, -10.25, -11.25, 12.25, 13.25, -14.25 }; + v (v2u8i, v4i16i, v8u32i, v16i64i, + v1flti, v2flti, v4flti, v8flti, + v1dbli, v2dbli, v4dbli, v8dbli); + } + +#if HAVE_INT128 + { + v1ldbl v1ldbli = { -0.125 }; + v2ldbl v2ldbli = { 1.125, -2.125 }; + v4ldbl v4ldbli = { -3.125, -4.125, 5.125, -6.125 }; + v8ldbl v8ldbli = { 7.125, -8.125, -9.125, 10.125, 11.125, 12.125, 13.125, 14.125 }; + vldbl (v1ldbli, v2ldbli, v4ldbli, v8ldbli); + } +#endif + + vla (argc); + + return 0; +} diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ef6397f1d28..61358ee5dcf 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,9 @@ 2019-06-18 Thomas Schwinge + * testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: New file. + * testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c: + Likewise. + * testsuite/libgomp.fortran/allocatable3.f90: Add missing results check. diff --git a/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C new file mode 100644 index 00000000000..c8dba9e5d1c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C @@ -0,0 +1,3 @@ +/* Verify OpenACC 'firstprivate' mappings for C++ reference types. */ + +#include "../../../gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c new file mode 100644 index 00000000000..4a8b310414c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c @@ -0,0 +1,6 @@ +/* Verify OpenACC 'firstprivate' mappings. */ + +/* { dg-additional-options "-Wno-psabi" } as apparently we're doing funny + things with vector arguments. */ + +#include "../../../gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c"