aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
diff options
context:
space:
mode:
Diffstat (limited to 'libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c')
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c125
1 files changed, 125 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
new file mode 100644
index 0000000..4b54171
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
@@ -0,0 +1,125 @@
+/* Inspired by 'gcc.target/nvptx/abi-struct-arg.c', 'gcc.target/nvptx/abi-struct-ret.c'. */
+
+/* See also '../libgomp.c-c++-common/target-abi-struct-1.c'. */
+
+/* To exercise PR119835 (if optimizations enabled): disable inlining, so that
+ GIMPLE passes still see the functions that return aggregate types. */
+#pragma GCC optimize "-fno-inline"
+
+typedef struct {} empty; /* See 'gcc/doc/extend.texi', "Empty Structures". */
+typedef struct {char a;} schar;
+typedef struct {short a;} sshort;
+typedef struct {int a;} sint;
+typedef struct {long long a;} slonglong;
+typedef struct {int a, b[12];} sint_13;
+
+#pragma omp declare target
+
+#define M(T) ({T t; t.a = sizeof t; t;})
+
+static __SIZE_TYPE__ empty_a;
+#pragma acc declare create(empty_a)
+#pragma acc routine
+static empty rempty(void)
+{
+ return ({empty t; empty_a = sizeof t; t;});
+}
+
+#pragma acc routine
+static schar rschar(void)
+{
+ return M(schar);
+}
+
+#pragma acc routine
+static sshort rsshort(void)
+{
+ return M(sshort);
+}
+
+#pragma acc routine
+static sint rsint(void)
+{
+ return M(sint);
+}
+
+#pragma acc routine
+static slonglong rslonglong(void)
+{
+ return M(slonglong);
+}
+
+#pragma acc routine
+static sint_13 rsint_13(void)
+{
+ return M(sint_13);
+}
+
+#pragma acc routine
+static void aempty(empty empty)
+{
+ (void) empty;
+
+ __SIZE_TYPE__ empty_a_exp;
+#ifndef __cplusplus
+ empty_a_exp = 0;
+#else
+ empty_a_exp = sizeof (char);
+#endif
+ if (empty_a != empty_a_exp)
+ __builtin_abort();
+}
+
+#pragma acc routine
+static void aschar(schar schar)
+{
+ if (schar.a != sizeof (char))
+ __builtin_abort();
+}
+
+#pragma acc routine
+static void asshort(sshort sshort)
+{
+ if (sshort.a != sizeof (short))
+ __builtin_abort();
+}
+
+#pragma acc routine
+static void asint(sint sint)
+{
+ if (sint.a != sizeof (int))
+ __builtin_abort();
+}
+
+#pragma acc routine
+static void aslonglong(slonglong slonglong)
+{
+ if (slonglong.a != sizeof (long long))
+ __builtin_abort();
+}
+
+#pragma acc routine
+static void asint_13(sint_13 sint_13)
+{
+ if (sint_13.a != (sizeof (int) * 13))
+ __builtin_abort();
+}
+
+#pragma omp end declare target
+
+int main()
+{
+#pragma omp target
+#pragma acc serial
+ /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */
+ {
+ aempty(rempty());
+ aschar(rschar());
+ asshort(rsshort());
+ asint(rsint());
+ aslonglong(rslonglong());
+ asint_13(rsint_13());
+ }
+
+ return 0;
+}