blob: 4b541711f36385e4628c80d0c3344076c80595fb (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
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;
}
|
