clang/test/SemaOpenCL/invalid-kernel-parameters.cl

   1 // RUN: %clang_cc1 -fsyntax-only -verify=expected,ocl12 %s -triple spir-unknown-unknown
   2 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown -cl-std=CL2.0
   3
   4 kernel 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
  13 kernel void no_ptrptr(global int * global *i) { }
  14 kernel void no_lptrcptr(constant int * local *i) { }
  15 kernel 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
  30 kernel void no_genericptr(generic int *ptr) { }
  31 // expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
  32 kernel 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}}
  34 kernel 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}}
  36 kernel 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
  40 void 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
  46 typedef __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}}
  48 typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'__private ptrdiff_t' (aka '__private int') declared here}}
  49 typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'__private intptr_t' (aka '__private int') declared here}}
  50 typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'__private uintptr_t' (aka '__private unsigned int') declared here}}
  51
  52 kernel 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
  54 kernel 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
  56 kernel 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
  58 kernel 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
  60 typedef size_t size_ty;
  61 struct 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 };
  64 kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'__private struct SizeTStruct' cannot be used as the type of a kernel parameter}}
  65
  66 union 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 };
  70 kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'__private union SizeTUnion' cannot be used as the type of a kernel parameter}}
  71
  72 typedef size_t s_ty; // expected-note{{'s_ty' (aka 'unsigned int') declared here}}
  73 typedef s_ty ss_ty; // expected-note{{'__private ss_ty' (aka '__private unsigned int') declared here}}
  74 kernel 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
  76 kernel 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.
  79 kernel void half_arg(half x) { }
  80
  81 typedef 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
  86 kernel 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
  90 typedef 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
  98 kernel void image_in_struct_arg(FooImage2D arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
  99
 100 typedef 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
 105 kernel void pointer_in_struct_arg(Foo arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
 106
 107 typedef struct FooGlobal // ocl12-note{{within field of type 'FooGlobal' declared here}}
 108 {
 109   global int* ptrField; // ocl12-note{{field of illegal pointer type '__global int *' declared here}}
 110 } FooGlobal;
 111
 112 kernel void global_pointer_in_struct_arg(FooGlobal arg) { } // ocl12-error{{struct kernel parameters may not contain pointers}}
 113
 114 typedef union FooUnion // expected-note{{within field of type 'FooUnion' declared here}}
 115 {
 116   int* ptrField; // expected-note-re{{field of illegal pointer type '__{{private|generic}} int *' declared here}}
 117 } FooUnion;
 118
 119 kernel void pointer_in_union_arg(FooUnion arg) { }// expected-error{{union kernel parameters may not contain pointers}}
 120
 121 typedef struct NestedPointer // expected-note 2 {{within field of type 'NestedPointer' declared here}}
 122 {
 123   int x;
 124   struct InnerNestedPointer
 125   {
 126     int* ptrField; // expected-note-re 3 {{field of illegal pointer type '__{{private|generic}} int *' declared here}}
 127   } inner; // expected-note 3 {{within field of type 'struct InnerNestedPointer' declared here}}
 128 } NestedPointer;
 129
 130 kernel void pointer_in_nested_struct_arg(NestedPointer arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 131
 132 struct NestedPointerComplex // expected-note{{within field of type 'NestedPointerComplex' declared here}}
 133 {
 134   int foo;
 135   float bar;
 136
 137   struct InnerNestedPointerComplex
 138   {
 139     int innerFoo;
 140     int* innerPtrField; // expected-note-re{{field of illegal pointer type '__{{private|generic}} int *' declared here}}
 141   } inner; // expected-note{{within field of type 'struct InnerNestedPointerComplex' declared here}}
 142
 143   float y;
 144   float z[4];
 145 };
 146
 147 kernel void pointer_in_nested_struct_arg_complex(struct NestedPointerComplex arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 148
 149 typedef struct NestedBool // expected-note 2 {{within field of type 'NestedBool' declared here}}
 150 {
 151   int x;
 152   struct InnerNestedBool
 153   {
 154     bool boolField; // expected-note 2 {{field of illegal type 'bool' declared here}}
 155   } inner; // expected-note 2 {{within field of type 'struct InnerNestedBool' declared here}}
 156 } NestedBool;
 157
 158 kernel 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}}
 159
 160 // Warning emitted again for argument used in other kernel
 161 kernel 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}}
 162
 163
 164 // Check for note with a struct not defined inside the struct
 165 typedef struct NestedBool2Inner
 166 {
 167   bool boolField; // expected-note{{field of illegal type 'bool' declared here}}
 168 } NestedBool2Inner;
 169
 170 typedef struct NestedBool2 // expected-note{{within field of type 'NestedBool2' declared here}}
 171 {
 172   int x;
 173   NestedBool2Inner inner; // expected-note{{within field of type 'NestedBool2Inner' (aka 'struct NestedBool2Inner') declared here}}
 174 } NestedBool2;
 175
 176 kernel 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}}
 177
 178
 179 struct InnerInner
 180 {
 181   int* foo;
 182   bool x;
 183 };
 184
 185 struct Valid
 186 {
 187   float c;
 188   float d;
 189 };
 190
 191 struct Inner
 192 {
 193   struct Valid v;
 194   struct InnerInner a;
 195   struct Valid g;
 196   struct InnerInner b;
 197 };
 198
 199 struct AlsoUser // expected-note{{within field of type 'AlsoUser' declared here}}
 200 {
 201   float x;
 202   struct Valid valid1;
 203   struct Valid valid2;
 204   struct NestedPointer aaaa; // expected-note{{within field of type 'struct NestedPointer' declared here}}
 205 };
 206
 207 kernel 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}}
 208
 209 struct ArrayOfPtr // expected-note{{within field of type 'ArrayOfPtr' declared here}}
 210 {
 211   float *arr[3]; // expected-note-re 2{{field of illegal type '__{{private|generic}} float *[3]' declared here}}
 212 };
 213 kernel void array_of_ptr(struct ArrayOfPtr arr) {} // expected-error{{struct kernel parameters may not contain pointers}}
 214
 215 struct ArrayOfStruct // expected-note{{within field of type 'ArrayOfStruct' declared here}}
 216 {
 217   struct ArrayOfPtr arr[3]; // expected-note{{within field of type 'struct ArrayOfPtr[3]' declared here}}
 218 };
 219 kernel void array_of_struct(struct ArrayOfStruct arr) {} // expected-error{{struct kernel parameters may not contain pointers}}