diff options
Diffstat (limited to 'test/SemaOpenCL')
| -rw-r--r-- | test/SemaOpenCL/access-qualifier.cl | 4 | ||||
| -rw-r--r-- | test/SemaOpenCL/address-spaces-conversions-cl2.0.cl | 37 | ||||
| -rw-r--r-- | test/SemaOpenCL/as_type.cl | 9 | ||||
| -rw-r--r-- | test/SemaOpenCL/atomic-init.cl | 12 | ||||
| -rw-r--r-- | test/SemaOpenCL/builtins-amdgcn-error-f16.cl | 6 | ||||
| -rw-r--r-- | test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl | 9 | ||||
| -rw-r--r-- | test/SemaOpenCL/builtins-amdgcn-error.cl | 50 | ||||
| -rw-r--r-- | test/SemaOpenCL/cl20-device-side-enqueue.cl | 21 | ||||
| -rw-r--r-- | test/SemaOpenCL/invalid-assignment-constant-address-space.cl | 7 | ||||
| -rw-r--r-- | test/SemaOpenCL/invalid-block.cl | 47 | ||||
| -rw-r--r-- | test/SemaOpenCL/invalid-logical-ops-1.1.cl | 57 | ||||
| -rw-r--r-- | test/SemaOpenCL/logical-ops.cl (renamed from test/SemaOpenCL/invalid-logical-ops-1.2.cl) | 59 | ||||
| -rw-r--r-- | test/SemaOpenCL/sampler_t.cl | 5 | ||||
| -rw-r--r-- | test/SemaOpenCL/types.cl | 6 | ||||
| -rw-r--r-- | test/SemaOpenCL/vector_swizzle_length.cl | 10 |
15 files changed, 247 insertions, 92 deletions
diff --git a/test/SemaOpenCL/access-qualifier.cl b/test/SemaOpenCL/access-qualifier.cl index 7bc974109f3c..35e838b1bf83 100644 --- a/test/SemaOpenCL/access-qualifier.cl +++ b/test/SemaOpenCL/access-qualifier.cl @@ -74,3 +74,7 @@ kernel void k14(read_only pipe int p) { myPipeWrite(p); // expected-error {{passing 'read_only pipe int' to parameter of incompatible type 'write_only pipe int'}} } #endif + +#if __OPENCL_C_VERSION__ < 200 +kernel void test_image3d_wo(write_only image3d_t img) {} // expected-error {{use of type '__write_only image3d_t' requires cl_khr_3d_image_writes extension to be enabled}} +#endif diff --git a/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl index 97fd07a24a24..73e29e7879d5 100644 --- a/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl +++ b/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl @@ -18,14 +18,20 @@ #ifdef GENERIC #define AS generic +#define AS_COMP local +#define AS_INCOMP constant #endif #ifdef GLOBAL #define AS global +#define AS_COMP global +#define AS_INCOMP local #endif #ifdef CONSTANT #define AS constant +#define AS_COMP constant +#define AS_INCOMP global #endif void f_glob(global int *arg_glob) {} @@ -263,12 +269,16 @@ void test_ternary() { var_void_gen = 0 ? var_cond : var_glob_ch; #ifdef CONSTANT // expected-error@-2{{conditional operator with the second and third operands of type ('__constant int *' and '__global char *') which are pointers to non-overlapping address spaces}} +#else +// expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__global char *')}} #endif local char *var_loc_ch; var_void_gen = 0 ? var_cond : var_loc_ch; #ifndef GENERIC // expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and '__local char *') which are pointers to non-overlapping address spaces}} +#else +// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__local char *')}} #endif constant void *var_void_const; @@ -276,18 +286,45 @@ void test_ternary() { var_void_const = 0 ? var_cond : var_const_ch; #ifndef CONSTANT // expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|generic}} int *' and '__constant char *') which are pointers to non-overlapping address spaces}} +#else +// expected-warning@-4{{pointer type mismatch ('__constant int *' and '__constant char *')}} #endif private char *var_priv_ch; var_void_gen = 0 ? var_cond : var_priv_ch; #ifndef GENERIC // expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}} +#else +// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}} #endif generic char *var_gen_ch; var_void_gen = 0 ? var_cond : var_gen_ch; #ifdef CONSTANT // expected-error@-2{{conditional operator with the second and third operands of type ('__constant int *' and '__generic char *') which are pointers to non-overlapping address spaces}} +#else +// expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}} #endif } +void test_pointer_chains() { + AS int *AS *var_as_as_int; + AS int *AS_COMP *var_asc_as_int; + AS_INCOMP int *AS_COMP *var_asc_asn_int; + AS_COMP int *AS_COMP *var_asc_asc_int; + + // Case 1: + // * address spaces of corresponded most outer pointees overlaps, their canonical types are equal + // * CVR, address spaces and canonical types of the rest of pointees are equivalent. + var_as_as_int = 0 ? var_as_as_int : var_asc_as_int; + + // Case 2: Corresponded inner pointees has non-overlapping address spaces. + var_as_as_int = 0 ? var_as_as_int : var_asc_asn_int; +// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}} + + // Case 3: Corresponded inner pointees has overlapping but not equivalent address spaces. +#ifdef GENERIC + var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int; +// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(local|global|constant)}} *')}} +#endif +} diff --git a/test/SemaOpenCL/as_type.cl b/test/SemaOpenCL/as_type.cl index f0bf4d7daef2..ad418097d902 100644 --- a/test/SemaOpenCL/as_type.cl +++ b/test/SemaOpenCL/as_type.cl @@ -1,7 +1,4 @@ -// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -verify -fsyntax-only - -typedef __attribute__(( ext_vector_type(3) )) char char3; -typedef __attribute__(( ext_vector_type(16) )) char char16; +// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -finclude-default-header -o - -verify -fsyntax-only char3 f1(char16 x) { return __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and 'char16' (vector of 16 'char' values) must match}} @@ -11,3 +8,7 @@ char16 f3(int x) { return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and 'int' must match}} } +void foo() { + char src = 1; + int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and 'char' must match}} +} diff --git a/test/SemaOpenCL/atomic-init.cl b/test/SemaOpenCL/atomic-init.cl new file mode 100644 index 000000000000..8208a85c3dab --- /dev/null +++ b/test/SemaOpenCL/atomic-init.cl @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -fsyntax-only -verify %s + +global atomic_int a1 = 0; + +kernel void test_atomic_initialization() { + a1 = 1; // expected-error {{atomic variable can be assigned to a variable only in global address space}} + atomic_int a2 = 0; // expected-error {{atomic variable can be initialized to a variable only in global address space}} + private atomic_int a3 = 0; // expected-error {{atomic variable can be initialized to a variable only in global address space}} + local atomic_int a4 = 0; // expected-error {{'__local' variable cannot have an initializer}} + global atomic_int a5 = 0; // expected-error {{function scope variable cannot be declared in global address space}} + static global atomic_int a6 = 0; +} diff --git a/test/SemaOpenCL/builtins-amdgcn-error-f16.cl b/test/SemaOpenCL/builtins-amdgcn-error-f16.cl index 7fa47179cb71..3487b1a5a803 100644 --- a/test/SemaOpenCL/builtins-amdgcn-error-f16.cl +++ b/test/SemaOpenCL/builtins-amdgcn-error-f16.cl @@ -1,9 +1,10 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -verify -S -o - %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable -void test_f16(global half *out, half a, half b, half c) +__attribute__((target("arch=tahiti"))) +void test_f16_tahiti(global half *out, half a, half b, half c) { *out = __builtin_amdgcn_div_fixuph(a, b, c); // expected-error {{'__builtin_amdgcn_div_fixuph' needs target feature 16-bit-insts}} *out = __builtin_amdgcn_rcph(a); // expected-error {{'__builtin_amdgcn_rcph' needs target feature 16-bit-insts}} @@ -15,4 +16,5 @@ void test_f16(global half *out, half a, half b, half c) *out = __builtin_amdgcn_frexp_exph(a); // expected-error {{'__builtin_amdgcn_frexp_exph' needs target feature 16-bit-insts}} *out = __builtin_amdgcn_fracth(a); // expected-error {{'__builtin_amdgcn_fracth' needs target feature 16-bit-insts}} *out = __builtin_amdgcn_classh(a, b); // expected-error {{'__builtin_amdgcn_classh' needs target feature 16-bit-insts}} + *out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}} } diff --git a/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl new file mode 100644 index 000000000000..c9fd8ab2cae8 --- /dev/null +++ b/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl @@ -0,0 +1,9 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu fiji -verify -S -o - %s + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +void test_gfx9_fmed3h(global half *out, half a, half b, half c) +{ + *out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}} +} diff --git a/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl index 83ccbefddc6f..2639bf27752f 100644 --- a/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -1,16 +1,17 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s -// FIXME: We only get one error if the functions are the other order in the -// file. - #pragma OPENCL EXTENSION cl_khr_fp64 : enable typedef unsigned long ulong; typedef unsigned int uint; -ulong test_s_memrealtime() +// To get all errors for feature checking we need to put them in one function +// since Clang will stop codegen for the next function if it finds error during +// codegen of the previous function. +void test_target_builtin(global int* out, int a) { - return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}} + __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, false); // expected-error {{'__builtin_amdgcn_mov_dpp' needs target feature dpp}} } void test_s_sleep(int x) @@ -18,6 +19,31 @@ void test_s_sleep(int x) __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}} } +void test_s_waitcnt(int x) +{ + __builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}} +} + +void test_s_sendmsg(int in) +{ + __builtin_amdgcn_s_sendmsg(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}} +} + +void test_s_sendmsg_var(int in1, int in2) +{ + __builtin_amdgcn_s_sendmsg(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}} +} + +void test_s_sendmsghalt(int in) +{ + __builtin_amdgcn_s_sendmsghalt(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}} +} + +void test_s_sendmsghalt_var(int in1, int in2) +{ + __builtin_amdgcn_s_sendmsghalt(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}} +} + void test_s_incperflevel(int x) { __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}} @@ -62,3 +88,17 @@ void test_ds_swizzle(global int* out, int a, int b) { *out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}} } + +void test_s_getreg(global int* out, int a) +{ + *out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}} +} + +void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e) +{ + *out = __builtin_amdgcn_mov_dpp(a, b, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, c, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, d, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} + *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}} +} + diff --git a/test/SemaOpenCL/cl20-device-side-enqueue.cl b/test/SemaOpenCL/cl20-device-side-enqueue.cl index c98145c2c900..3e87cfcddf66 100644 --- a/test/SemaOpenCL/cl20-device-side-enqueue.cl +++ b/test/SemaOpenCL/cl20-device-side-enqueue.cl @@ -1,6 +1,7 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -triple "spir-unknown-unknown" -verify -pedantic -fsyntax-only -DB32 // RUN: %clang_cc1 %s -cl-std=CL2.0 -triple "spir64-unknown-unknown" -verify -pedantic -fsyntax-only -Wconversion -DWCONV +typedef struct {int a;} ndrange_t; // Diagnostic tests for different overloads of enqueue_kernel from Table 6.13.17.1 of OpenCL 2.0 Spec. kernel void enqueue_kernel_tests() { queue_t default_queue; @@ -30,7 +31,7 @@ kernel void enqueue_kernel_tests() { enqueue_kernel(default_queue, flags, ndrange, vptr); // expected-error{{illegal call to enqueue_kernel, expected block argument}} - enqueue_kernel(default_queue, flags, ndrange, ^(int i) { // expected-error{{blocks in this form of device side enqueue call are expected to have have no parameters}} + enqueue_kernel(default_queue, flags, ndrange, ^(int i) { // expected-error{{blocks with parameters are not accepted in this prototype of enqueue_kernel call}} return 0; }); @@ -111,7 +112,7 @@ kernel void enqueue_kernel_tests() { const bl_B_t block_B = (bl_B_t) ^ (local void *a, local int *b) {}; - enqueue_kernel(default_queue, flags, ndrange, block_B, 1024, 1024); // expected-error{{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + enqueue_kernel(default_queue, flags, ndrange, block_B, 1024, 1024); // expected-error{{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} enqueue_kernel(default_queue, flags, ndrange, // expected-error{{mismatch in number of block parameters and local size arguments passed}} ^(local void *a, local void *b) { @@ -177,12 +178,12 @@ kernel void work_group_size_tests() { size = get_kernel_work_group_size(^(local void *a) { return; }); - size = get_kernel_work_group_size(^(local int *a) { // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + size = get_kernel_work_group_size(^(local int *a) { // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} return; }); - size = get_kernel_work_group_size(block_B); // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} - size = get_kernel_work_group_size(block_D); // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} - size = get_kernel_work_group_size(^(int a) { // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + size = get_kernel_work_group_size(block_B); // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} + size = get_kernel_work_group_size(block_D); // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} + size = get_kernel_work_group_size(^(int a) { // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} return; }); size = get_kernel_work_group_size(); // expected-error {{too few arguments to function call, expected 1, have 0}} @@ -194,14 +195,14 @@ kernel void work_group_size_tests() { size = get_kernel_preferred_work_group_size_multiple(^(local void *a) { return; }); - size = get_kernel_preferred_work_group_size_multiple(^(local int *a) { // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + size = get_kernel_preferred_work_group_size_multiple(^(local int *a) { // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} return; }); - size = get_kernel_preferred_work_group_size_multiple(^(int a) { // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + size = get_kernel_preferred_work_group_size_multiple(^(int a) { // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} return; }); - size = get_kernel_preferred_work_group_size_multiple(block_B); // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} - size = get_kernel_preferred_work_group_size_multiple(block_D); // expected-error {{blocks used in device side enqueue are expected to have parameters of type 'local void*'}} + size = get_kernel_preferred_work_group_size_multiple(block_B); // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} + size = get_kernel_preferred_work_group_size_multiple(block_D); // expected-error {{blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'}} size = get_kernel_preferred_work_group_size_multiple(); // expected-error {{too few arguments to function call, expected 1, have 0}} size = get_kernel_preferred_work_group_size_multiple(1); // expected-error{{expected block argument}} size = get_kernel_preferred_work_group_size_multiple(block_A, 1); // expected-error{{too many arguments to function call, expected 1, have 2}} diff --git a/test/SemaOpenCL/invalid-assignment-constant-address-space.cl b/test/SemaOpenCL/invalid-assignment-constant-address-space.cl new file mode 100644 index 000000000000..32917c493ce3 --- /dev/null +++ b/test/SemaOpenCL/invalid-assignment-constant-address-space.cl @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only + +int constant c[3] = {0}; + +void foo() { + c[0] = 1; //expected-error{{read-only variable is not assignable}} +} diff --git a/test/SemaOpenCL/invalid-block.cl b/test/SemaOpenCL/invalid-block.cl index a9eac1c2b582..89bf03264e13 100644 --- a/test/SemaOpenCL/invalid-block.cl +++ b/test/SemaOpenCL/invalid-block.cl @@ -4,26 +4,34 @@ void f0(int (^const bl)()); // All blocks declarations must be const qualified and initialized. void f1() { - int (^bl1)() = ^() {return 1;}; - int (^const bl2)() = ^(){return 1;}; + int (^bl1)(void) = ^() { + return 1; + }; + int (^const bl2)(void) = ^() { + return 1; + }; f0(bl1); f0(bl2); - bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (^const)()' and 'int (^const)()')}} + bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (__generic ^const)(void)' and 'int (__generic ^const)(void)')}} int (^const bl3)(); // expected-error{{invalid block variable declaration - must be initialized}} } // A block with extern storage class is not allowed. -extern int (^bl)() = ^(){return 1;}; // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}} +extern int (^bl)(void) = ^() { // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}} + return 1; +}; void f2() { - extern int (^bl)() = ^(){return 1;}; // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}} + extern int (^bl)(void) = ^() { // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}} + return 1; + }; } // A block cannot be the return value of a function. typedef int (^bl_t)(void); -bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (^const)(void)') is not allowed}} +bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (__generic ^const)(void)') is not allowed}} struct bl_s { - int (^bl)(void); // expected-error {{the 'int (^const)(void)' type cannot be used to declare a structure or union field}} + int (^bl)(void); // expected-error {{the 'int (__generic ^const)(void)' type cannot be used to declare a structure or union field}} }; void f4() { @@ -45,16 +53,31 @@ void f5(int i) { bl2_t bl2 = ^(int i) { return 2; }; - bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (^const)(int)') type is invalid in OpenCL}} + bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(int)') type is invalid in OpenCL}} int tmp = i ? bl1(i) // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}} : bl2(i); // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}} } // A block pointer type and all pointer operations are disallowed -void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (^const __generic)(int)') is invalid in OpenCL}} +void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}} bl2_t bl = ^(int i) { return 1; }; - bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (^const __generic)(int)') is invalid in OpenCL}} - *bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}} - &bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}} + bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}} + *bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}} + &bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}} +} +// A block can't reference another block +kernel void f7() { + bl2_t bl1 = ^(int i) { + return 1; + }; + void (^bl2)(void) = ^{ + int i = bl1(1); // expected-error {{cannot refer to a block inside block}} + }; + void (^bl3)(void) = ^{ + }; + void (^bl4)(void) = ^{ + bl3(); // expected-error {{cannot refer to a block inside block}} + }; + return; } diff --git a/test/SemaOpenCL/invalid-logical-ops-1.1.cl b/test/SemaOpenCL/invalid-logical-ops-1.1.cl deleted file mode 100644 index 2269dd322b36..000000000000 --- a/test/SemaOpenCL/invalid-logical-ops-1.1.cl +++ /dev/null @@ -1,57 +0,0 @@ -// RUN: %clang_cc1 %s -verify -cl-std=CL1.1 -triple x86_64-unknown-linux-gnu - -#pragma OPENCL EXTENSION cl_khr_fp64 : enable -typedef __attribute__((ext_vector_type(4))) float float4; -typedef __attribute__((ext_vector_type(4))) double double4; -typedef __attribute__((ext_vector_type(4))) int int4; -typedef __attribute__((ext_vector_type(4))) long long4; - -kernel void float_ops() { - int flaf = 0.0f && 0.0f; // expected-error {{invalid operands}} - int flof = 0.0f || 0.0f; // expected-error {{invalid operands}} - float fbaf = 0.0f & 0.0f; // expected-error {{invalid operands}} - float fbof = 0.0f | 0.0f; // expected-error {{invalid operands}} - float fbxf = 0.0f ^ 0.0f; // expected-error {{invalid operands}} - int flai = 0.0f && 0; // expected-error {{invalid operands}} - int floi = 0.0f || 0; // expected-error {{invalid operands}} - float ibaf = 0 & 0.0f; // expected-error {{invalid operands}} - float ibof = 0 | 0.0f; // expected-error {{invalid operands}} - float bnf = ~0.0f; // expected-error {{invalid argument type}} - float lnf = !0.0f; // expected-error {{invalid argument type}} -} - -kernel void vec_float_ops() { - float4 f4 = (float4)(0, 0, 0, 0); - int4 f4laf = f4 && 0.0f; // expected-error {{invalid operands}} - int4 f4lof = f4 || 0.0f; // expected-error {{invalid operands}} - float4 f4baf = f4 & 0.0f; // expected-error {{invalid operands}} - float4 f4bof = f4 | 0.0f; // expected-error {{invalid operands}} - float4 f4bxf = f4 ^ 0.0f; // expected-error {{invalid operands}} - float bnf4 = ~f4; // expected-error {{invalid argument type}} - int4 lnf4 = !f4; // expected-error {{invalid argument type}} -} - -kernel void double_ops() { - int flaf = 0.0 && 0.0; // expected-error {{invalid operands}} - int flof = 0.0 || 0.0; // expected-error {{invalid operands}} - double fbaf = 0.0 & 0.0; // expected-error {{invalid operands}} - double fbof = 0.0 | 0.0; // expected-error {{invalid operands}} - double fbxf = 0.0 ^ 0.0; // expected-error {{invalid operands}} - int flai = 0.0 && 0; // expected-error {{invalid operands}} - int floi = 0.0 || 0; // expected-error {{invalid operands}} - double ibaf = 0 & 0.0; // expected-error {{invalid operands}} - double ibof = 0 | 0.0; // expected-error {{invalid operands}} - double bnf = ~0.0; // expected-error {{invalid argument type}} - double lnf = !0.0; // expected-error {{invalid argument type}} -} - -kernel void vec_double_ops() { - double4 f4 = (double4)(0, 0, 0, 0); - long4 f4laf = f4 && 0.0; // expected-error {{invalid operands}} - long4 f4lof = f4 || 0.0; // expected-error {{invalid operands}} - double4 f4baf = f4 & 0.0; // expected-error {{invalid operands}} - double4 f4bof = f4 | 0.0; // expected-error {{invalid operands}} - double4 f4bxf = f4 ^ 0.0; // expected-error {{invalid operands}} - double bnf4 = ~f4; // expected-error {{invalid argument type}} - long4 lnf4 = !f4; // expected-error {{invalid argument type}} -} diff --git a/test/SemaOpenCL/invalid-logical-ops-1.2.cl b/test/SemaOpenCL/logical-ops.cl index bee52396cc6c..42501b14413e 100644 --- a/test/SemaOpenCL/invalid-logical-ops-1.2.cl +++ b/test/SemaOpenCL/logical-ops.cl @@ -1,4 +1,6 @@ +// RUN: %clang_cc1 %s -verify -cl-std=CL1.1 -triple x86_64-unknown-linux-gnu // RUN: %clang_cc1 %s -verify -cl-std=CL1.2 -triple x86_64-unknown-linux-gnu +// RUN: %clang_cc1 %s -verify -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu #pragma OPENCL EXTENSION cl_khr_fp64 : enable @@ -9,50 +11,107 @@ typedef __attribute__((ext_vector_type(4))) long long4; kernel void float_ops() { int flaf = 0.0f && 0.0f; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif int flof = 0.0f || 0.0f; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif float fbaf = 0.0f & 0.0f; // expected-error {{invalid operands}} float fbof = 0.0f | 0.0f; // expected-error {{invalid operands}} float fbxf = 0.0f ^ 0.0f; // expected-error {{invalid operands}} int flai = 0.0f && 0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif int floi = 0.0f || 0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif float ibaf = 0 & 0.0f; // expected-error {{invalid operands}} float ibof = 0 | 0.0f; // expected-error {{invalid operands}} float bnf = ~0.0f;// expected-error {{invalid argument type}} float lnf = !0.0f; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid argument type}} +#endif } kernel void vec_float_ops() { float4 f4 = (float4)(0, 0, 0, 0); int4 f4laf = f4 && 0.0f; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif int4 f4lof = f4 || 0.0f; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif float4 f4baf = f4 & 0.0f; // expected-error {{invalid operands}} float4 f4bof = f4 | 0.0f; // expected-error {{invalid operands}} float4 f4bxf = f4 ^ 0.0f; // expected-error {{invalid operands}} float bnf4 = ~f4; // expected-error {{invalid argument type}} int4 lnf4 = !f4; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid argument type}} +#endif } kernel void double_ops() { int flaf = 0.0 && 0.0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif int flof = 0.0 || 0.0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif double fbaf = 0.0 & 0.0; // expected-error {{invalid operands}} double fbof = 0.0 | 0.0; // expected-error {{invalid operands}} double fbxf = 0.0 ^ 0.0; // expected-error {{invalid operands}} int flai = 0.0 && 0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif int floi = 0.0 || 0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif double ibaf = 0 & 0.0; // expected-error {{invalid operands}} double ibof = 0 | 0.0; // expected-error {{invalid operands}} double bnf = ~0.0; // expected-error {{invalid argument type}} double lnf = !0.0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid argument type}} +#endif } kernel void vec_double_ops() { double4 f4 = (double4)(0, 0, 0, 0); long4 f4laf = f4 && 0.0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif long4 f4lof = f4 || 0.0; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid operands}} +#endif double4 f4baf = f4 & 0.0; // expected-error {{invalid operands}} double4 f4bof = f4 | 0.0; // expected-error {{invalid operands}} double4 f4bxf = f4 ^ 0.0; // expected-error {{invalid operands}} double bnf4 = ~f4; // expected-error {{invalid argument type}} long4 lnf4 = !f4; +#if __OPENCL_C_VERSION__ < 120 +// expected-error@-2{{invalid argument type}} +#endif +} + +kernel void pointer_ops(){ + global int* p; + bool b = !p; + b = p==0; + int i; + b = !&i; + b = &i==(int *)1; } diff --git a/test/SemaOpenCL/sampler_t.cl b/test/SemaOpenCL/sampler_t.cl index c87b6da7c7e1..0dddeeb39013 100644 --- a/test/SemaOpenCL/sampler_t.cl +++ b/test/SemaOpenCL/sampler_t.cl @@ -30,7 +30,7 @@ constant sampler_t glb_smp8 = 1.0f; // expected-error{{initializing '__constant constant sampler_t glb_smp9 = 0x100000000LL; // expected-error{{sampler_t initialization requires 32-bit integer, not 'long long'}} -void foo(sampler_t); +void foo(sampler_t); // expected-note{{passing argument to parameter here}} constant struct sampler_s { sampler_t smp; // expected-error{{the 'sampler_t' type cannot be used to declare a structure or union field}} @@ -65,7 +65,8 @@ void kernel ker(sampler_t argsmp) { foo(const_smp5); foo(const_smp6); foo(argsmp); - foo(5); // expected-error{{sampler_t variable required - got 'int'}} + foo(5); + foo(5.0f); // expected-error {{passing 'float' to parameter of incompatible type 'sampler_t'}} sampler_t sa[] = {argsmp, const_smp}; // expected-error {{array of 'sampler_t' type is invalid in OpenCL}} foo(sa[0]); foo(bad()); diff --git a/test/SemaOpenCL/types.cl b/test/SemaOpenCL/types.cl new file mode 100644 index 000000000000..dc14800f3532 --- /dev/null +++ b/test/SemaOpenCL/types.cl @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only + +// expected-no-diagnostics + +// Check redefinition of standard types +typedef atomic_int atomic_flag; diff --git a/test/SemaOpenCL/vector_swizzle_length.cl b/test/SemaOpenCL/vector_swizzle_length.cl new file mode 100644 index 000000000000..94e3f654d5d9 --- /dev/null +++ b/test/SemaOpenCL/vector_swizzle_length.cl @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only + +typedef float float8 __attribute__((ext_vector_type(8))); + +void foo() { + float8 f2 = (float8)(0, 0, 0, 0, 0, 0, 0, 0); + + f2.s01234; // expected-error {{vector component access has invalid length 5. Supported: 1,2,3,4,8,16}} + f2.xyzxy; // expected-error {{vector component access has invalid length 5. Supported: 1,2,3,4,8,16}} +} |