1// RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
2// RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown -cl-std=CL2.0
3
4kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
5
6#pragma OPENCL EXTENSION cl_khr_fp16 : enable
7
8#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
9// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
10// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
11// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
12#endif
13kernel void no_ptrptr(global int * global *i) { }
14kernel void no_lptrcptr(constant int * local *i) { }
15kernel void no_ptrptrptr(global int * global * global *i) { }
16
17// expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
18__kernel void no_privateptr(__private int *i) { }
19
20// expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
21__kernel void no_privatearray(__private int i[]) { }
22
23// expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
24__kernel void no_addrsp_ptr(int *ptr) { }
25
26// expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
27__kernel void no_defaultarray(int i[]) { }
28
29#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
30kernel void no_genericptr(generic int *ptr) { }
31// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
32kernel void no_ptr_private_ptr(private int * global *i) { }
33// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
34kernel void no_ptr_ptr_private_ptr(private int * global * global *i) { }
35// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
36kernel void no_ptr_private_ptr_ptr(global int * private * global *i) { }
37// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
38#endif
39
40void no_addrspace_param(global int x) { } // expected-error{{parameter may not be qualified with an address space}}
41
42// Disallowed: parameters with type
43// bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t
44// or a struct / union with any of these types in them
45
46typedef __SIZE_TYPE__ size_t; // expected-note{{'__private size_t' (aka '__private unsigned int') declared here}}
47                              // expected-note@-1{{'size_t' (aka 'unsigned int') declared here}}
48typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'__private ptrdiff_t' (aka '__private int') declared here}}
49typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'__private intptr_t' (aka '__private int') declared here}}
50typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'__private uintptr_t' (aka '__private unsigned int') declared here}}
51
52kernel void size_t_arg(size_t x) {} // expected-error{{'__private size_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
53
54kernel void ptrdiff_t_arg(ptrdiff_t x) {} // expected-error{{'__private ptrdiff_t' (aka '__private int') cannot be used as the type of a kernel parameter}}
55
56kernel void intptr_t_arg(intptr_t x) {} // expected-error{{'__private intptr_t' (aka '__private int') cannot be used as the type of a kernel parameter}}
57
58kernel void uintptr_t_arg(uintptr_t x) {} // expected-error{{'__private uintptr_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
59
60typedef size_t size_ty;
61struct SizeTStruct { // expected-note{{within field of type 'SizeTStruct' declared here}}
62  size_ty s; // expected-note{{field of illegal type 'size_ty' (aka 'unsigned int') declared here}}
63};
64kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'__private struct SizeTStruct' cannot be used as the type of a kernel parameter}}
65
66union SizeTUnion { // expected-note{{within field of type 'SizeTUnion' declared here}}
67  size_t s; // expected-note{{field of illegal type 'size_t' (aka 'unsigned int') declared here}}
68  float f;
69};
70kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'__private union SizeTUnion' cannot be used as the type of a kernel parameter}}
71
72typedef size_t s_ty; // expected-note{{'s_ty' (aka 'unsigned int') declared here}}
73typedef s_ty ss_ty; // expected-note{{'__private ss_ty' (aka '__private unsigned int') declared here}}
74kernel void typedef_to_size_t(ss_ty s) {} // expected-error{{'__private ss_ty' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
75
76kernel void bool_arg(bool x) { } // expected-error{{'__private bool' cannot be used as the type of a kernel parameter}}
77
78// half kernel argument is allowed when cl_khr_fp16 is enabled.
79kernel void half_arg(half x) { }
80
81typedef struct ContainsBool // expected-note{{within field of type 'ContainsBool' declared here}}
82{
83  bool x; // expected-note{{field of illegal type 'bool' declared here}}
84} ContainsBool;
85
86kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'__private ContainsBool' (aka '__private struct ContainsBool') cannot be used as the type of a kernel parameter}}
87
88
89
90typedef struct FooImage2D // expected-note{{within field of type 'FooImage2D' declared here}}
91{
92  // TODO: Clean up needed - we don't really need to check for image, event, etc
93  // as a note here any longer.
94  // They are diagnosed as an error for all struct fields (OpenCL v1.2 s6.9b,r).
95  image2d_t imageField; // expected-note{{field of illegal type '__read_only image2d_t' declared here}} expected-error{{the '__read_only image2d_t' type cannot be used to declare a structure or union field}}
96} FooImage2D;
97
98kernel void image_in_struct_arg(FooImage2D arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
99
100typedef struct Foo // expected-note{{within field of type 'Foo' declared here}}
101{
102  int* ptrField; // expected-note-re{{field of illegal pointer type '__{{private|generic}} int *' declared here}}
103} Foo;
104
105kernel void pointer_in_struct_arg(Foo arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
106
107typedef union FooUnion // expected-note{{within field of type 'FooUnion' declared here}}
108{
109  int* ptrField; // expected-note-re{{field of illegal pointer type '__{{private|generic}} int *' declared here}}
110} FooUnion;
111
112kernel void pointer_in_union_arg(FooUnion arg) { }// expected-error{{union kernel parameters may not contain pointers}}
113
114typedef struct NestedPointer // expected-note 2 {{within field of type 'NestedPointer' declared here}}
115{
116  int x;
117  struct InnerNestedPointer
118  {
119    int* ptrField; // expected-note-re 3 {{field of illegal pointer type '__{{private|generic}} int *' declared here}}
120  } inner; // expected-note 3 {{within field of type 'struct InnerNestedPointer' declared here}}
121} NestedPointer;
122
123kernel void pointer_in_nested_struct_arg(NestedPointer arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
124
125struct NestedPointerComplex // expected-note{{within field of type 'NestedPointerComplex' declared here}}
126{
127  int foo;
128  float bar;
129
130  struct InnerNestedPointerComplex
131  {
132    int innerFoo;
133    int* innerPtrField; // expected-note-re{{field of illegal pointer type '__{{private|generic}} int *' declared here}}
134  } inner; // expected-note{{within field of type 'struct InnerNestedPointerComplex' declared here}}
135
136  float y;
137  float z[4];
138};
139
140kernel void pointer_in_nested_struct_arg_complex(struct NestedPointerComplex arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
141
142typedef struct NestedBool // expected-note 2 {{within field of type 'NestedBool' declared here}}
143{
144  int x;
145  struct InnerNestedBool
146  {
147    bool boolField; // expected-note 2 {{field of illegal type 'bool' declared here}}
148  } inner; // expected-note 2 {{within field of type 'struct InnerNestedBool' declared here}}
149} NestedBool;
150
151kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
152
153// Warning emitted again for argument used in other kernel
154kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
155
156
157// Check for note with a struct not defined inside the struct
158typedef struct NestedBool2Inner
159{
160  bool boolField; // expected-note{{field of illegal type 'bool' declared here}}
161} NestedBool2Inner;
162
163typedef struct NestedBool2 // expected-note{{within field of type 'NestedBool2' declared here}}
164{
165  int x;
166  NestedBool2Inner inner; // expected-note{{within field of type 'NestedBool2Inner' (aka 'struct NestedBool2Inner') declared here}}
167} NestedBool2;
168
169kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'__private NestedBool2' (aka '__private struct NestedBool2') cannot be used as the type of a kernel parameter}}
170
171
172struct InnerInner
173{
174  int* foo;
175  bool x;
176};
177
178struct Valid
179{
180  float c;
181  float d;
182};
183
184struct Inner
185{
186  struct Valid v;
187  struct InnerInner a;
188  struct Valid g;
189  struct InnerInner b;
190};
191
192struct AlsoUser // expected-note{{within field of type 'AlsoUser' declared here}}
193{
194  float x;
195  struct Valid valid1;
196  struct Valid valid2;
197  struct NestedPointer aaaa; // expected-note{{within field of type 'struct NestedPointer' declared here}}
198};
199
200kernel void pointer_in_nested_struct_arg_2(struct Valid valid, struct NestedPointer arg, struct AlsoUser also) { } // expected-error 2 {{struct kernel parameters may not contain pointers}}
201
202struct ArrayOfPtr // expected-note{{within field of type 'ArrayOfPtr' declared here}}
203{
204  float *arr[3]; // expected-note-re 2{{field of illegal type '__{{private|generic}} float *[3]' declared here}}
205};
206kernel void array_of_ptr(struct ArrayOfPtr arr) {} // expected-error{{struct kernel parameters may not contain pointers}}
207
208struct ArrayOfStruct // expected-note{{within field of type 'ArrayOfStruct' declared here}}
209{
210  struct ArrayOfPtr arr[3]; // expected-note{{within field of type 'struct ArrayOfPtr [3]' declared here}}
211};
212kernel void array_of_struct(struct ArrayOfStruct arr) {} // expected-error{{struct kernel parameters may not contain pointers}}
213