diff options
Diffstat (limited to 'test')
40 files changed, 795 insertions, 277 deletions
diff --git a/test/CodeGen/builtins-ppc-error.c b/test/CodeGen/builtins-ppc-error.c new file mode 100644 index 000000000000..5860c4f9e77e --- /dev/null +++ b/test/CodeGen/builtins-ppc-error.c @@ -0,0 +1,20 @@ +// REQUIRES: powerpc-registered-target + +// RUN: %clang_cc1 -faltivec -target-feature +power9-vector \ +// RUN: -triple powerpc64-unknown-unknown -fsyntax-only \ +// RUN: -Wall -Werror -verify %s + +// RUN: %clang_cc1 -faltivec -target-feature +power9-vector \ +// RUN: -triple powerpc64le-unknown-unknown -fsyntax-only \ +// RUN: -Wall -Werror -verify %s + +#include <altivec.h> + +extern vector signed int vsi; +extern vector unsigned char vuc; + +void testInsertWord1(void) { + int index = 5; + vector unsigned char v1 = vec_insert4b(vsi, vuc, index); // expected-error {{argument to '__builtin_vsx_insertword' must be a constant integer}} + vector unsigned long long v2 = vec_extract4b(vuc, index); // expected-error {{argument to '__builtin_vsx_extractuword' must be a constant integer}} +} diff --git a/test/CodeGen/builtins-ppc-p9vector.c b/test/CodeGen/builtins-ppc-p9vector.c index f70d2f9f1504..bd0ad182f15f 100644 --- a/test/CodeGen/builtins-ppc-p9vector.c +++ b/test/CodeGen/builtins-ppc-p9vector.c @@ -1166,17 +1166,52 @@ vector float test114(void) { // CHECK-BE: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 undef, i32 0, i32 undef, i32 1, i32 undef, i32 2, i32 undef, i32 3> // CHECK-BE: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) // CHECK-BE-NEXT: ret <4 x float> -// CHECK-LE: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 0, i32 undef, i32 1, i32 undef, i32 2, i32 undef, i32 3, i32 undef> -// CHECK-LE: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) -// CHECK-LE-NEXT: ret <4 x float> +// CHECK: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 0, i32 undef, i32 1, i32 undef, i32 2, i32 undef, i32 3, i32 undef> +// CHECK: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) +// CHECK-NEXT: ret <4 x float> return vec_extract_fp32_from_shorth(vusa); } vector float test115(void) { // CHECK-BE: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 undef, i32 4, i32 undef, i32 5, i32 undef, i32 6, i32 undef, i32 7> // CHECK-BE: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) // CHECK-BE-NEXT: ret <4 x float> -// CHECK-LE: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 4, i32 undef, i32 5, i32 undef, i32 6, i32 undef, i32 7, i32 undef> -// CHECK-LE: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) -// CHECK-LE-NEXT: ret <4 x float> +// CHECK: shufflevector <8 x i16> {{.+}}, <8 x i16> {{.+}}, <8 x i32> <i32 4, i32 undef, i32 5, i32 undef, i32 6, i32 undef, i32 7, i32 undef> +// CHECK: @llvm.ppc.vsx.xvcvhpsp(<8 x i16> {{.+}}) +// CHECK-NEXT: ret <4 x float> return vec_extract_fp32_from_shortl(vusa); } +vector unsigned char test116(void) { +// CHECK-BE: [[T1:%.+]] = call <4 x i32> @llvm.ppc.vsx.xxinsertw(<4 x i32> {{.+}}, <2 x i64> {{.+}}, i32 7) +// CHECK-BE-NEXT: bitcast <4 x i32> [[T1]] to <16 x i8> +// CHECK: [[T1:%.+]] = shufflevector <2 x i64> {{.+}}, <2 x i64> {{.+}}, <2 x i32> <i32 1, i32 0> +// CHECK-NEXT: [[T2:%.+]] = bitcast <2 x i64> [[T1]] to <4 x i32> +// CHECK-NEXT: [[T3:%.+]] = call <4 x i32> @llvm.ppc.vsx.xxinsertw(<4 x i32> [[T2]], <2 x i64> {{.+}}, i32 5) +// CHECK-NEXT: bitcast <4 x i32> [[T3]] to <16 x i8> + return vec_insert4b(vuia, vuca, 7); +} +vector unsigned char test117(void) { +// CHECK-BE: [[T1:%.+]] = call <4 x i32> @llvm.ppc.vsx.xxinsertw(<4 x i32> {{.+}}, <2 x i64> {{.+}}, i32 12) +// CHECK-BE-NEXT: bitcast <4 x i32> [[T1]] to <16 x i8> +// CHECK: [[T1:%.+]] = shufflevector <2 x i64> {{.+}}, <2 x i64> {{.+}}, <2 x i32> <i32 1, i32 0> +// CHECK-NEXT: [[T2:%.+]] = bitcast <2 x i64> [[T1]] to <4 x i32> +// CHECK-NEXT: [[T3:%.+]] = call <4 x i32> @llvm.ppc.vsx.xxinsertw(<4 x i32> [[T2]], <2 x i64> {{.+}}, i32 0) +// CHECK-NEXT: bitcast <4 x i32> [[T3]] to <16 x i8> + return vec_insert4b(vuia, vuca, 13); +} +vector unsigned long long test118(void) { +// CHECK-BE: call <2 x i64> @llvm.ppc.vsx.xxextractuw(<2 x i64> {{.+}}, i32 11) +// CHECK-BE-NEXT: ret <2 x i64> +// CHECK: [[T1:%.+]] = call <2 x i64> @llvm.ppc.vsx.xxextractuw(<2 x i64> {{.+}}, i32 1) +// CHECK-NEXT: shufflevector <2 x i64> [[T1]], <2 x i64> [[T1]], <2 x i32> <i32 1, i32 0> +// CHECK-NEXT: ret <2 x i64> + return vec_extract4b(vuca, 11); +} +vector unsigned long long test119(void) { +// CHECK-BE: call <2 x i64> @llvm.ppc.vsx.xxextractuw(<2 x i64> {{.+}}, i32 0) +// CHECK-BE-NEXT: ret <2 x i64> +// CHECK: [[T1:%.+]] = call <2 x i64> @llvm.ppc.vsx.xxextractuw(<2 x i64> {{.+}}, i32 12) +// CHECK-NEXT: shufflevector <2 x i64> [[T1]], <2 x i64> [[T1]], <2 x i32> <i32 1, i32 0> +// CHECK-NEXT: ret <2 x i64> + return vec_extract4b(vuca, -5); +} + diff --git a/test/CodeGen/catch-undef-behavior.c b/test/CodeGen/catch-undef-behavior.c index c2f01ae1a66f..d7a26f8a7d4b 100644 --- a/test/CodeGen/catch-undef-behavior.c +++ b/test/CodeGen/catch-undef-behavior.c @@ -6,16 +6,16 @@ // CHECK-UBSAN: @[[INT:.*]] = private unnamed_addr constant { i16, i16, [6 x i8] } { i16 0, i16 11, [6 x i8] c"'int'\00" } // FIXME: When we only emit each type once, use [[INT]] more below. -// CHECK-UBSAN: @[[LINE_100:.*]] = private unnamed_addr global {{.*}}, i32 100, i32 5 {{.*}} @[[INT]], i64 4, i8 1 -// CHECK-UBSAN: @[[LINE_200:.*]] = {{.*}}, i32 200, i32 10 {{.*}}, i64 4, i8 0 +// CHECK-UBSAN: @[[LINE_100:.*]] = private unnamed_addr global {{.*}}, i32 100, i32 5 {{.*}} @[[INT]], i8 2, i8 1 +// CHECK-UBSAN: @[[LINE_200:.*]] = {{.*}}, i32 200, i32 10 {{.*}}, i8 2, i8 0 // CHECK-UBSAN: @[[LINE_300:.*]] = {{.*}}, i32 300, i32 12 {{.*}} @{{.*}}, {{.*}} @{{.*}} // CHECK-UBSAN: @[[LINE_400:.*]] = {{.*}}, i32 400, i32 12 {{.*}} @{{.*}}, {{.*}} @{{.*}} -// CHECK-UBSAN: @[[LINE_500:.*]] = {{.*}}, i32 500, i32 10 {{.*}} @{{.*}}, i64 4, i8 0 } -// CHECK-UBSAN: @[[LINE_600:.*]] = {{.*}}, i32 600, i32 3 {{.*}} @{{.*}}, i64 4, i8 1 } +// CHECK-UBSAN: @[[LINE_500:.*]] = {{.*}}, i32 500, i32 10 {{.*}} @{{.*}}, i8 2, i8 0 } +// CHECK-UBSAN: @[[LINE_600:.*]] = {{.*}}, i32 600, i32 3 {{.*}} @{{.*}}, i8 2, i8 1 } // CHECK-UBSAN: @[[STRUCT_S:.*]] = private unnamed_addr constant { i16, i16, [11 x i8] } { i16 -1, i16 0, [11 x i8] c"'struct S'\00" } -// CHECK-UBSAN: @[[LINE_700:.*]] = {{.*}}, i32 700, i32 14 {{.*}} @[[STRUCT_S]], i64 4, i8 3 } +// CHECK-UBSAN: @[[LINE_700:.*]] = {{.*}}, i32 700, i32 14 {{.*}} @[[STRUCT_S]], i8 2, i8 3 } // CHECK-UBSAN: @[[LINE_800:.*]] = {{.*}}, i32 800, i32 12 {{.*}} @{{.*}} } // CHECK-UBSAN: @[[LINE_900:.*]] = {{.*}}, i32 900, i32 11 {{.*}} @{{.*}} } // CHECK-UBSAN: @[[LINE_1000:.*]] = {{.*}}, i32 1000, i32 10 {{.*}} @{{.*}} } @@ -54,7 +54,7 @@ void foo() { // CHECK-TRAP: br i1 %[[OK]], {{.*}} // CHECK-UBSAN: %[[ARG:.*]] = ptrtoint {{.*}} %[[PTR]] to i64 - // CHECK-UBSAN-NEXT: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_100]] to i8*), i64 %[[ARG]]) + // CHECK-UBSAN-NEXT: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_100]] to i8*), i64 %[[ARG]]) // CHECK-TRAP: call void @llvm.trap() [[NR_NUW:#[0-9]+]] // CHECK-TRAP-NEXT: unreachable @@ -62,7 +62,7 @@ void foo() { // With -fsanitize=null, only perform the null check. // CHECK-NULL: %[[NULL:.*]] = icmp ne {{.*}}, null // CHECK-NULL: br i1 %[[NULL]] - // CHECK-NULL: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_100]] to i8*), i64 %{{.*}}) + // CHECK-NULL: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_100]] to i8*), i64 %{{.*}}) #line 100 u.i=1; } @@ -77,7 +77,7 @@ int bar(int *a) { // CHECK-COMMON-NEXT: icmp eq i64 %[[MISALIGN]], 0 // CHECK-UBSAN: %[[ARG:.*]] = ptrtoint - // CHECK-UBSAN-NEXT: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_200]] to i8*), i64 %[[ARG]]) + // CHECK-UBSAN-NEXT: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_200]] to i8*), i64 %[[ARG]]) // CHECK-TRAP: call void @llvm.trap() [[NR_NUW]] // CHECK-TRAP-NEXT: unreachable @@ -145,7 +145,7 @@ int rsh_inbounds(int a, int b) { // CHECK-COMMON-LABEL: @load int load(int *p) { - // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_500]] to i8*), i64 %{{.*}}) + // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_500]] to i8*), i64 %{{.*}}) // CHECK-TRAP: call void @llvm.trap() [[NR_NUW]] // CHECK-TRAP-NEXT: unreachable @@ -155,7 +155,7 @@ int load(int *p) { // CHECK-COMMON-LABEL: @store void store(int *p, int q) { - // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_600]] to i8*), i64 %{{.*}}) + // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_600]] to i8*), i64 %{{.*}}) // CHECK-TRAP: call void @llvm.trap() [[NR_NUW]] // CHECK-TRAP-NEXT: unreachable @@ -167,7 +167,7 @@ struct S { int k; }; // CHECK-COMMON-LABEL: @member_access int *member_access(struct S *p) { - // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch(i8* bitcast ({{.*}} @[[LINE_700]] to i8*), i64 %{{.*}}) + // CHECK-UBSAN: call void @__ubsan_handle_type_mismatch_v1(i8* bitcast ({{.*}} @[[LINE_700]] to i8*), i64 %{{.*}}) // CHECK-TRAP: call void @llvm.trap() [[NR_NUW]] // CHECK-TRAP-NEXT: unreachable diff --git a/test/CodeGen/sanitize-recover.c b/test/CodeGen/sanitize-recover.c index b263f5163181..dd8734e971eb 100644 --- a/test/CodeGen/sanitize-recover.c +++ b/test/CodeGen/sanitize-recover.c @@ -33,7 +33,7 @@ void foo() { // PARTIAL: br i1 %[[CHECK012]], {{.*}} !prof ![[WEIGHT_MD:.*]], !nosanitize // PARTIAL: br i1 %[[CHECK02]], {{.*}} - // PARTIAL: call void @__ubsan_handle_type_mismatch_abort( + // PARTIAL: call void @__ubsan_handle_type_mismatch_v1_abort( // PARTIAL-NEXT: unreachable - // PARTIAL: call void @__ubsan_handle_type_mismatch( + // PARTIAL: call void @__ubsan_handle_type_mismatch_v1( } diff --git a/test/CodeGen/vectorcall.c b/test/CodeGen/vectorcall.c index b38d5e5fbc5b..167f72ca2cfd 100644 --- a/test/CodeGen/vectorcall.c +++ b/test/CodeGen/vectorcall.c @@ -1,22 +1,22 @@ -// RUN: %clang_cc1 -emit-llvm %s -o - -triple=i386-pc-win32 | FileCheck %s -// RUN: %clang_cc1 -emit-llvm %s -o - -triple=x86_64-pc-win32 | FileCheck %s --check-prefix=X64 +// RUN: %clang_cc1 -emit-llvm %s -o - -ffreestanding -triple=i386-pc-win32 | FileCheck %s --check-prefix=X32 +// RUN: %clang_cc1 -emit-llvm %s -o - -ffreestanding -triple=x86_64-pc-win32 | FileCheck %s --check-prefix=X64 void __vectorcall v1(int a, int b) {} -// CHECK: define x86_vectorcallcc void @"\01v1@@8"(i32 inreg %a, i32 inreg %b) +// X32: define x86_vectorcallcc void @"\01v1@@8"(i32 inreg %a, i32 inreg %b) // X64: define x86_vectorcallcc void @"\01v1@@16"(i32 %a, i32 %b) void __vectorcall v2(char a, char b) {} -// CHECK: define x86_vectorcallcc void @"\01v2@@8"(i8 inreg signext %a, i8 inreg signext %b) +// X32: define x86_vectorcallcc void @"\01v2@@8"(i8 inreg signext %a, i8 inreg signext %b) // X64: define x86_vectorcallcc void @"\01v2@@16"(i8 %a, i8 %b) struct Small { int x; }; void __vectorcall v3(int a, struct Small b, int c) {} -// CHECK: define x86_vectorcallcc void @"\01v3@@12"(i32 inreg %a, i32 %b.0, i32 inreg %c) +// X32: define x86_vectorcallcc void @"\01v3@@12"(i32 inreg %a, i32 %b.0, i32 inreg %c) // X64: define x86_vectorcallcc void @"\01v3@@24"(i32 %a, i32 %b.coerce, i32 %c) struct Large { int a[5]; }; void __vectorcall v4(int a, struct Large b, int c) {} -// CHECK: define x86_vectorcallcc void @"\01v4@@28"(i32 inreg %a, %struct.Large* byval align 4 %b, i32 inreg %c) +// X32: define x86_vectorcallcc void @"\01v4@@28"(i32 inreg %a, %struct.Large* byval align 4 %b, i32 inreg %c) // X64: define x86_vectorcallcc void @"\01v4@@40"(i32 %a, %struct.Large* %b, i32 %c) struct HFA2 { double x, y; }; @@ -24,54 +24,84 @@ struct HFA4 { double w, x, y, z; }; struct HFA5 { double v, w, x, y, z; }; void __vectorcall hfa1(int a, struct HFA4 b, int c) {} -// CHECK: define x86_vectorcallcc void @"\01hfa1@@40"(i32 inreg %a, double %b.0, double %b.1, double %b.2, double %b.3, i32 inreg %c) -// X64: define x86_vectorcallcc void @"\01hfa1@@48"(i32 %a, double %b.0, double %b.1, double %b.2, double %b.3, i32 %c) +// X32: define x86_vectorcallcc void @"\01hfa1@@40"(i32 inreg %a, %struct.HFA4 inreg %b.coerce, i32 inreg %c) +// X64: define x86_vectorcallcc void @"\01hfa1@@48"(i32 %a, %struct.HFA4 inreg %b.coerce, i32 %c) // HFAs that would require more than six total SSE registers are passed // indirectly. Additional vector arguments can consume the rest of the SSE // registers. void __vectorcall hfa2(struct HFA4 a, struct HFA4 b, double c) {} -// CHECK: define x86_vectorcallcc void @"\01hfa2@@72"(double %a.0, double %a.1, double %a.2, double %a.3, %struct.HFA4* inreg %b, double %c) -// X64: define x86_vectorcallcc void @"\01hfa2@@72"(double %a.0, double %a.1, double %a.2, double %a.3, %struct.HFA4* %b, double %c) +// X32: define x86_vectorcallcc void @"\01hfa2@@72"(%struct.HFA4 inreg %a.coerce, %struct.HFA4* inreg %b, double %c) +// X64: define x86_vectorcallcc void @"\01hfa2@@72"(%struct.HFA4 inreg %a.coerce, %struct.HFA4* %b, double %c) // Ensure that we pass builtin types directly while counting them against the // SSE register usage. void __vectorcall hfa3(double a, double b, double c, double d, double e, struct HFA2 f) {} -// CHECK: define x86_vectorcallcc void @"\01hfa3@@56"(double %a, double %b, double %c, double %d, double %e, %struct.HFA2* inreg %f) +// X32: define x86_vectorcallcc void @"\01hfa3@@56"(double %a, double %b, double %c, double %d, double %e, %struct.HFA2* inreg %f) // X64: define x86_vectorcallcc void @"\01hfa3@@56"(double %a, double %b, double %c, double %d, double %e, %struct.HFA2* %f) // Aggregates with more than four elements are not HFAs and are passed byval. // Because they are not classified as homogeneous, they don't get special // handling to ensure alignment. void __vectorcall hfa4(struct HFA5 a) {} -// CHECK: define x86_vectorcallcc void @"\01hfa4@@40"(%struct.HFA5* byval align 4) +// X32: define x86_vectorcallcc void @"\01hfa4@@40"(%struct.HFA5* byval align 4) // X64: define x86_vectorcallcc void @"\01hfa4@@40"(%struct.HFA5* %a) // Return HFAs of 4 or fewer elements in registers. static struct HFA2 g_hfa2; struct HFA2 __vectorcall hfa5(void) { return g_hfa2; } -// CHECK: define x86_vectorcallcc %struct.HFA2 @"\01hfa5@@0"() +// X32: define x86_vectorcallcc %struct.HFA2 @"\01hfa5@@0"() // X64: define x86_vectorcallcc %struct.HFA2 @"\01hfa5@@0"() typedef float __attribute__((vector_size(16))) v4f32; struct HVA2 { v4f32 x, y; }; +struct HVA3 { v4f32 w, x, y; }; struct HVA4 { v4f32 w, x, y, z; }; +struct HVA5 { v4f32 w, x, y, z, p; }; -void __vectorcall hva1(int a, struct HVA4 b, int c) {} -// CHECK: define x86_vectorcallcc void @"\01hva1@@72"(i32 inreg %a, <4 x float> %b.0, <4 x float> %b.1, <4 x float> %b.2, <4 x float> %b.3, i32 inreg %c) -// X64: define x86_vectorcallcc void @"\01hva1@@80"(i32 %a, <4 x float> %b.0, <4 x float> %b.1, <4 x float> %b.2, <4 x float> %b.3, i32 %c) +v4f32 __vectorcall hva1(int a, struct HVA4 b, int c) {return b.w;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva1@@72"(i32 inreg %a, %struct.HVA4 inreg %b.coerce, i32 inreg %c) +// X64: define x86_vectorcallcc <4 x float> @"\01hva1@@80"(i32 %a, %struct.HVA4 inreg %b.coerce, i32 %c) -void __vectorcall hva2(struct HVA4 a, struct HVA4 b, v4f32 c) {} -// CHECK: define x86_vectorcallcc void @"\01hva2@@144"(<4 x float> %a.0, <4 x float> %a.1, <4 x float> %a.2, <4 x float> %a.3, %struct.HVA4* inreg %b, <4 x float> %c) -// X64: define x86_vectorcallcc void @"\01hva2@@144"(<4 x float> %a.0, <4 x float> %a.1, <4 x float> %a.2, <4 x float> %a.3, %struct.HVA4* %b, <4 x float> %c) +v4f32 __vectorcall hva2(struct HVA4 a, struct HVA4 b, v4f32 c) {return c;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva2@@144"(%struct.HVA4 inreg %a.coerce, %struct.HVA4* inreg %b, <4 x float> %c) +// X64: define x86_vectorcallcc <4 x float> @"\01hva2@@144"(%struct.HVA4 inreg %a.coerce, %struct.HVA4* %b, <4 x float> %c) -void __vectorcall hva3(v4f32 a, v4f32 b, v4f32 c, v4f32 d, v4f32 e, struct HVA2 f) {} -// CHECK: define x86_vectorcallcc void @"\01hva3@@112"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, <4 x float> %e, %struct.HVA2* inreg %f) -// X64: define x86_vectorcallcc void @"\01hva3@@112"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, <4 x float> %e, %struct.HVA2* %f) +v4f32 __vectorcall hva3(v4f32 a, v4f32 b, v4f32 c, v4f32 d, v4f32 e, struct HVA2 f) {return f.x;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva3@@112"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, <4 x float> %e, %struct.HVA2* inreg %f) +// X64: define x86_vectorcallcc <4 x float> @"\01hva3@@112"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, <4 x float> %e, %struct.HVA2* %f) + +// vector types have higher priority then HVA structures, So vector types are allocated first +// and HVAs are allocated if enough registers are available +v4f32 __vectorcall hva4(struct HVA4 a, struct HVA2 b, v4f32 c) {return b.y;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva4@@112"(%struct.HVA4 inreg %a.coerce, %struct.HVA2* inreg %b, <4 x float> %c) +// X64: define x86_vectorcallcc <4 x float> @"\01hva4@@112"(%struct.HVA4 inreg %a.coerce, %struct.HVA2* %b, <4 x float> %c) + +v4f32 __vectorcall hva5(struct HVA3 a, struct HVA3 b, v4f32 c, struct HVA2 d) {return d.y;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva5@@144"(%struct.HVA3 inreg %a.coerce, %struct.HVA3* inreg %b, <4 x float> %c, %struct.HVA2 inreg %d.coerce) +// X64: define x86_vectorcallcc <4 x float> @"\01hva5@@144"(%struct.HVA3 inreg %a.coerce, %struct.HVA3* %b, <4 x float> %c, %struct.HVA2 inreg %d.coerce) + +struct HVA4 __vectorcall hva6(struct HVA4 a, struct HVA4 b) { return b;} +// X32: define x86_vectorcallcc %struct.HVA4 @"\01hva6@@128"(%struct.HVA4 inreg %a.coerce, %struct.HVA4* inreg %b) +// X64: define x86_vectorcallcc %struct.HVA4 @"\01hva6@@128"(%struct.HVA4 inreg %a.coerce, %struct.HVA4* %b) + +struct HVA5 __vectorcall hva7() {struct HVA5 a = {}; return a;} +// X32: define x86_vectorcallcc void @"\01hva7@@0"(%struct.HVA5* inreg noalias sret %agg.result) +// X64: define x86_vectorcallcc void @"\01hva7@@0"(%struct.HVA5* noalias sret %agg.result) + +v4f32 __vectorcall hva8(v4f32 a, v4f32 b, v4f32 c, v4f32 d, int e, v4f32 f) {return f;} +// X32: define x86_vectorcallcc <4 x float> @"\01hva8@@84"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, i32 inreg %e, <4 x float> %f) +// X64: define x86_vectorcallcc <4 x float> @"\01hva8@@88"(<4 x float> %a, <4 x float> %b, <4 x float> %c, <4 x float> %d, i32 %e, <4 x float> %f) typedef float __attribute__((ext_vector_type(3))) v3f32; struct OddSizeHVA { v3f32 x, y; }; void __vectorcall odd_size_hva(struct OddSizeHVA a) {} -// CHECK: define x86_vectorcallcc void @"\01odd_size_hva@@32"(<3 x float> %a.0, <3 x float> %a.1) -// X64: define x86_vectorcallcc void @"\01odd_size_hva@@32"(<3 x float> %a.0, <3 x float> %a.1) +// X32: define x86_vectorcallcc void @"\01odd_size_hva@@32"(%struct.OddSizeHVA inreg %a.coerce) +// X64: define x86_vectorcallcc void @"\01odd_size_hva@@32"(%struct.OddSizeHVA inreg %a.coerce) + +// The Vectorcall ABI only allows passing the first 6 items in registers, so this shouldn't +// consider 'p7' as a register. Instead p5 gets put into the register on the second pass. +struct HFA2 __vectorcall AddParticles(struct HFA2 p1, float p2, struct HFA4 p3, int p4, struct HFA2 p5, float p6, float p7){ return p1;} +// X32: define x86_vectorcallcc %struct.HFA2 @"\01AddParticles@@80"(%struct.HFA2 inreg %p1.coerce, float %p2, %struct.HFA4* inreg %p3, i32 inreg %p4, %struct.HFA2 inreg %p5.coerce, float %p6, float %p7) +// X64: define x86_vectorcallcc %struct.HFA2 @"\01AddParticles@@96"(%struct.HFA2 inreg %p1.coerce, float %p2, %struct.HFA4* %p3, i32 %p4, %struct.HFA2 inreg %p5.coerce, float %p6, float %p7) diff --git a/test/CodeGenCXX/dllexport.cpp b/test/CodeGenCXX/dllexport.cpp index eb9ca79b7b40..116176e2cb92 100644 --- a/test/CodeGenCXX/dllexport.cpp +++ b/test/CodeGenCXX/dllexport.cpp @@ -515,6 +515,18 @@ struct __declspec(dllexport) ClassWithClosure { // M32-DAG: ret void }; +template <typename T> struct TemplateWithClosure { + TemplateWithClosure(int x = sizeof(T)) {} +}; +extern template struct TemplateWithClosure<char>; +template struct __declspec(dllexport) TemplateWithClosure<char>; +extern template struct TemplateWithClosure<int>; +template struct __declspec(dllexport) TemplateWithClosure<int>; +// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01??_F?$TemplateWithClosure@D@@QAEXXZ"({{.*}}) {{#[0-9]+}} comdat +// M32-DAG: call {{.*}} @"\01??0?$TemplateWithClosure@D@@QAE@H@Z"({{.*}}, i32 1) +// M32-DAG: define weak_odr dllexport x86_thiscallcc void @"\01??_F?$TemplateWithClosure@H@@QAEXXZ"({{.*}}) {{#[0-9]+}} comdat +// M32-DAG: call {{.*}} @"\01??0?$TemplateWithClosure@H@@QAE@H@Z"({{.*}}, i32 4) + struct __declspec(dllexport) NestedOuter { DELETE_IMPLICIT_MEMBERS(NestedOuter); NestedOuter(void *p = 0) {} diff --git a/test/CodeGenCXX/homogeneous-aggregates.cpp b/test/CodeGenCXX/homogeneous-aggregates.cpp index 67911c0d7f90..1338b25e21ae 100644 --- a/test/CodeGenCXX/homogeneous-aggregates.cpp +++ b/test/CodeGenCXX/homogeneous-aggregates.cpp @@ -47,7 +47,7 @@ D1 CC func_D1(D1 x) { return x; } // PPC: define [3 x double] @_Z7func_D22D2([3 x double] %x.coerce) // ARM32: define arm_aapcs_vfpcc %struct.D2 @_Z7func_D22D2(%struct.D2 %x.coerce) // ARM64: define %struct.D2 @_Z7func_D22D2([3 x double] %x.coerce) -// X64: define x86_vectorcallcc %struct.D2 @"\01_Z7func_D22D2@@24"(double %x.0, double %x.1, double %x.2) +// X64: define x86_vectorcallcc %struct.D2 @"\01_Z7func_D22D2@@24"(%struct.D2 inreg %x.coerce) D2 CC func_D2(D2 x) { return x; } // PPC: define void @_Z7func_D32D3(%struct.D3* noalias sret %agg.result, [4 x i64] %x.coerce) @@ -92,7 +92,7 @@ struct HVAWithEmptyBase : Float1, Empty, Float2 { float z; }; void CC with_empty_base(HVAWithEmptyBase a) {} // FIXME: MSVC doesn't consider this an HVA because of the empty base. -// X64: define x86_vectorcallcc void @"\01_Z15with_empty_base16HVAWithEmptyBase@@16"(float %a.0, float %a.1, float %a.2) +// X64: define x86_vectorcallcc void @"\01_Z15with_empty_base16HVAWithEmptyBase@@16"(%struct.HVAWithEmptyBase inreg %a.coerce) struct HVAWithEmptyBitField : Float1, Float2 { int : 0; // Takes no space. @@ -102,5 +102,5 @@ struct HVAWithEmptyBitField : Float1, Float2 { // PPC: define void @_Z19with_empty_bitfield20HVAWithEmptyBitField([3 x float] %a.coerce) // ARM64: define void @_Z19with_empty_bitfield20HVAWithEmptyBitField([3 x float] %a.coerce) // ARM32: define arm_aapcs_vfpcc void @_Z19with_empty_bitfield20HVAWithEmptyBitField(%struct.HVAWithEmptyBitField %a.coerce) -// X64: define x86_vectorcallcc void @"\01_Z19with_empty_bitfield20HVAWithEmptyBitField@@16"(float %a.0, float %a.1, float %a.2) +// X64: define x86_vectorcallcc void @"\01_Z19with_empty_bitfield20HVAWithEmptyBitField@@16"(%struct.HVAWithEmptyBitField inreg %a.coerce) void CC with_empty_bitfield(HVAWithEmptyBitField a) {} diff --git a/test/CodeGenCXX/ubsan-vtable-checks.cpp b/test/CodeGenCXX/ubsan-vtable-checks.cpp index 80af77d4ea6d..e684ae9180f1 100644 --- a/test/CodeGenCXX/ubsan-vtable-checks.cpp +++ b/test/CodeGenCXX/ubsan-vtable-checks.cpp @@ -21,7 +21,7 @@ int get_v(T* t) { // CHECK-NULL-NOT: load {{.*}} (%struct.T*{{.*}})**, {{.*}} (%struct.T*{{.*}})*** // CHECK-NULL: [[UBSAN_CMP_RES:%[0-9]+]] = icmp ne %struct.T* %{{[_a-z0-9]+}}, null // CHECK-NULL-NEXT: br i1 [[UBSAN_CMP_RES]], label %{{.*}}, label %{{.*}} - // CHECK-NULL: call void @__ubsan_handle_type_mismatch_abort + // CHECK-NULL: call void @__ubsan_handle_type_mismatch_v1_abort // Second, we check that vtable is actually loaded once the type check is done. // CHECK-NULL: load {{.*}} (%struct.T*{{.*}})**, {{.*}} (%struct.T*{{.*}})*** return t->v(); diff --git a/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/bin/.keep b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/bin/.keep new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/bin/.keep diff --git a/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/include/.keep b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/include/.keep new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/include/.keep diff --git a/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/lib/.keep b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/lib/.keep new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/lib/.keep diff --git a/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_30.10.bc b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_30.10.bc new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_30.10.bc diff --git a/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_35.10.bc b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_35.10.bc new file mode 100644 index 000000000000..e69de29bb2d1 --- /dev/null +++ b/test/Driver/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/nvvm/libdevice/libdevice.compute_35.10.bc diff --git a/test/Driver/avr-toolchain.c b/test/Driver/avr-toolchain.c new file mode 100644 index 000000000000..46a3c10fa3a1 --- /dev/null +++ b/test/Driver/avr-toolchain.c @@ -0,0 +1,4 @@ +// A basic clang -cc1 command-line. + +// RUN: %clang %s -### -no-canonical-prefixes -target avr 2>&1 | FileCheck -check-prefix=CC1 %s +// CC1: clang{{.*}} "-cc1" "-triple" "avr" diff --git a/test/Driver/cuda-version-check.cu b/test/Driver/cuda-version-check.cu index cb2ac7994f75..46ca72f2ea0c 100644 --- a/test/Driver/cuda-version-check.cu +++ b/test/Driver/cuda-version-check.cu @@ -2,40 +2,40 @@ // REQUIRES: x86-registered-target // REQUIRES: nvptx-registered-target -// RUN: %clang -v -### --cuda-gpu-arch=sm_20 --sysroot=%S/Inputs/CUDA 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_20 --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=OK -// RUN: %clang -v -### --cuda-gpu-arch=sm_20 --sysroot=%S/Inputs/CUDA_80 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_20 --sysroot=%S/Inputs/CUDA_80 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=OK -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA_80 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA_80 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=OK // The installation at Inputs/CUDA is CUDA 7.0, which doesn't support sm_60. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=ERR_SM60 // This should only complain about sm_60, not sm_35. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --cuda-gpu-arch=sm_35 \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --cuda-gpu-arch=sm_35 \ // RUN: --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=ERR_SM60 --check-prefix=OK_SM35 // We should get two errors here, one for sm_60 and one for sm_61. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --cuda-gpu-arch=sm_61 \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --cuda-gpu-arch=sm_61 \ // RUN: --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=ERR_SM60 --check-prefix=ERR_SM61 // We should still get an error if we pass -nocudainc, because this compilation // would invoke ptxas, and we do a version check on that, too. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 -nocudainc --sysroot=%S/Inputs/CUDA 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 -nocudainc --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=ERR_SM60 // If with -nocudainc and -E, we don't touch the CUDA install, so we // shouldn't get an error. -// RUN: %clang -v -### -E --cuda-device-only --cuda-gpu-arch=sm_60 -nocudainc \ +// RUN: %clang --target=x86_64-linux -v -### -E --cuda-device-only --cuda-gpu-arch=sm_60 -nocudainc \ // RUN: --sysroot=%S/Inputs/CUDA 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=OK // --no-cuda-version-check should suppress all of these errors. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA 2>&1 \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --sysroot=%S/Inputs/CUDA 2>&1 \ // RUN: --no-cuda-version-check %s | \ // RUN: FileCheck %s --check-prefix=OK @@ -43,9 +43,9 @@ // therefore we should not get an error in host-only mode. We use the -S here // to avoid the error being produced in case by the assembler tool, which does // the same check. -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --cuda-host-only --sysroot=%S/Inputs/CUDA -S 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --cuda-host-only --sysroot=%S/Inputs/CUDA -S 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=OK -// RUN: %clang -v -### --cuda-gpu-arch=sm_60 --cuda-device-only --sysroot=%S/Inputs/CUDA -S 2>&1 %s | \ +// RUN: %clang --target=x86_64-linux -v -### --cuda-gpu-arch=sm_60 --cuda-device-only --sysroot=%S/Inputs/CUDA -S 2>&1 %s | \ // RUN: FileCheck %s --check-prefix=ERR_SM60 // OK-NOT: error: GPU arch diff --git a/test/Driver/cuda-windows.cu b/test/Driver/cuda-windows.cu new file mode 100644 index 000000000000..1d67710647c0 --- /dev/null +++ b/test/Driver/cuda-windows.cu @@ -0,0 +1,14 @@ +// REQUIRES: clang-driver +// REQUIRES: x86-registered-target +// REQUIRES: nvptx-registered-target +// +// RUN: %clang -v --target=i386-pc-windows-msvc \ +// RUN: --sysroot=%S/Inputs/CUDA-windows 2>&1 %s -### | FileCheck %s +// RUN: %clang -v --target=i386-pc-windows-mingw32 \ +// RUN: --sysroot=%S/Inputs/CUDA-windows 2>&1 %s -### | FileCheck %s + +// CHECK: Found CUDA installation: {{.*}}/Inputs/CUDA-windows/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0 +// CHECK: "-cc1" "-triple" "nvptx-nvidia-cuda" +// CHECK-SAME: "-fms-extensions" +// CHECK-SAME: "-fms-compatibility" +// CHECK-SAME: "-fms-compatibility-version= diff --git a/test/Index/complete-block-properties.m b/test/Index/complete-block-properties.m index d166147294e1..4697703c8e5c 100644 --- a/test/Index/complete-block-properties.m +++ b/test/Index/complete-block-properties.m @@ -43,7 +43,7 @@ typedef int (^BarBlock)(int *); //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void}{TypedText block}{LeftParen (}{RightParen )} (35) //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void (^)()}{TypedText block}{Equal = }{Placeholder ^(void)} (38) //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType Foo}{TypedText blocker}{LeftParen (}{Placeholder int x}{Comma , }{Placeholder Foo y}{Comma , }{Placeholder ^(Foo *someParameter)foo}{RightParen )} (35) -//CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType Foo (^)(int, Foo, FooBlock)}{TypedText blocker}{Equal = }{Placeholder ^Foo(int x, Foo y, FooBlock foo)} (38) +//CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType Foo (^)(int, Foo, FooBlock)}{TypedText blocker}{Equal = }{Placeholder ^Foo(int x, Foo y, FooBlock foo)} (32) //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType int}{TypedText foo} (35) //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void}{TypedText fooBlock}{LeftParen (}{Placeholder Foo *someParameter}{RightParen )} (35) //CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType Test *}{TypedText getObject}{LeftParen (}{Placeholder int index}{RightParen )} (35) diff --git a/test/Index/complete-block-property-assignment.m b/test/Index/complete-block-property-assignment.m index ced3b7fa1302..908e18629528 100644 --- a/test/Index/complete-block-property-assignment.m +++ b/test/Index/complete-block-property-assignment.m @@ -15,6 +15,7 @@ typedef void (^FooBlock)(Foo *someParameter); @interface Test : Obj @property (readwrite, nonatomic, copy) FooBlock onEventHandler; @property (readonly, nonatomic, copy) void (^onReadonly)(int *someParameter); +@property (readwrite, nonatomic, copy) int (^processEvent)(int eventCode); @property (readonly, nonatomic, strong) Obj *obj; @end @@ -29,10 +30,10 @@ typedef void (^FooBlock)(Foo *someParameter); SELFY.foo = 2 } -// RUN: c-index-test -code-completion-at=%s:26:8 %s | FileCheck -check-prefix=CHECK-CC1 %s -// RUN: c-index-test -code-completion-at=%s:27:27 %s | FileCheck -check-prefix=CHECK-CC1 %s -// RUN: c-index-test -code-completion-at=%s:28:22 %s | FileCheck -check-prefix=CHECK-CC1 %s -// RUN: c-index-test -code-completion-at=%s:29:9 %s | FileCheck -check-prefix=CHECK-CC1 %s +// RUN: c-index-test -code-completion-at=%s:27:8 %s | FileCheck -check-prefix=CHECK-CC1 %s +// RUN: c-index-test -code-completion-at=%s:28:27 %s | FileCheck -check-prefix=CHECK-CC1 %s +// RUN: c-index-test -code-completion-at=%s:29:22 %s | FileCheck -check-prefix=CHECK-CC1 %s +// RUN: c-index-test -code-completion-at=%s:30:9 %s | FileCheck -check-prefix=CHECK-CC1 %s // CHECK-CC1: ObjCPropertyDecl:{ResultType int}{TypedText foo} (35) // CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType Obj *}{TypedText obj} (35) // CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void}{TypedText onAction}{LeftParen (}{Placeholder Obj *object}{RightParen )} (35) @@ -40,6 +41,8 @@ typedef void (^FooBlock)(Foo *someParameter); // CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void}{TypedText onEventHandler}{LeftParen (}{Placeholder Foo *someParameter}{RightParen )} (35) // CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType FooBlock}{TypedText onEventHandler}{Equal = }{Placeholder ^(Foo *someParameter)} (38) // CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType void}{TypedText onReadonly}{LeftParen (}{Placeholder int *someParameter}{RightParen )} (35) +// CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType int}{TypedText processEvent}{LeftParen (}{Placeholder int eventCode}{RightParen )} (35) +// CHECK-CC1-NEXT: ObjCPropertyDecl:{ResultType int (^)(int)}{TypedText processEvent}{Equal = }{Placeholder ^int(int eventCode)} (32) - (void) takeInt:(int)x { } @@ -53,16 +56,17 @@ typedef void (^FooBlock)(Foo *someParameter); return self.foo; } -// RUN: c-index-test -code-completion-at=%s:47:9 %s | FileCheck -check-prefix=CHECK-NO %s -// RUN: c-index-test -code-completion-at=%s:48:16 %s | FileCheck -check-prefix=CHECK-NO %s -// RUN: c-index-test -code-completion-at=%s:49:23 %s | FileCheck -check-prefix=CHECK-NO %s -// RUN: c-index-test -code-completion-at=%s:50:12 %s | FileCheck -check-prefix=CHECK-NO %s -// RUN: c-index-test -code-completion-at=%s:51:15 %s | FileCheck -check-prefix=CHECK-NO %s -// RUN: c-index-test -code-completion-at=%s:53:15 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:50:9 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:51:16 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:52:23 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:53:12 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:54:15 %s | FileCheck -check-prefix=CHECK-NO %s +// RUN: c-index-test -code-completion-at=%s:56:15 %s | FileCheck -check-prefix=CHECK-NO %s // CHECK-NO: ObjCPropertyDecl:{ResultType int}{TypedText foo} (35) // CHECK-NO-NEXT: ObjCPropertyDecl:{ResultType Obj *}{TypedText obj} (35) // CHECK-NO-NEXT: ObjCPropertyDecl:{ResultType void (^)(Obj *)}{TypedText onAction} (35) // CHECK-NO-NEXT: ObjCPropertyDecl:{ResultType FooBlock}{TypedText onEventHandler} (35) // CHECK-NO-NEXT: ObjCPropertyDecl:{ResultType void (^)(int *)}{TypedText onReadonly} (35) +// CHECK-NO-NEXT: ObjCPropertyDecl:{ResultType int (^)(int)}{TypedText processEvent} (35) @end diff --git a/test/OpenMP/nvptx_target_codegen.cpp b/test/OpenMP/nvptx_target_codegen.cpp index 287089d7c45e..59c4d5b277ce 100644 --- a/test/OpenMP/nvptx_target_codegen.cpp +++ b/test/OpenMP/nvptx_target_codegen.cpp @@ -8,9 +8,6 @@ #ifndef HEADER #define HEADER -// CHECK-DAG: [[OMP_NT:@.+]] = common addrspace(3) global i32 0 -// CHECK-DAG: [[OMP_WID:@.+]] = common addrspace(3) global i64 0 - template<typename tx, typename ty> struct TT{ tx X; @@ -26,19 +23,22 @@ int foo(int n) { double cn[5][n]; TT<long long, char> d; - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l87}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l90}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -54,31 +54,34 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l87]]() - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l90]]() + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] // CHECK: {{call|invoke}} void [[T1]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: br label {{%?}}[[EXIT:.+]] // - // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) - // CHECK: br label {{%?}}[[TERM:.+]] + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[MASTER]] + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] + // CHECK: br label {{%?}}[[TERMINATE:.+]] + // + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // @@ -93,19 +96,22 @@ int foo(int n) { { } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l158}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l167}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -121,35 +127,38 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l158]](i[[SZ:32|64]] [[ARG1:%[^)]+]]) + // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l167]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]]) // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], // CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]], // CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16* - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] - // CHECK: {{call|invoke}} void [[T3]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: {{call|invoke}} void [[T2]]_worker() + // CHECK: br label {{%?}}[[EXIT:.+]] + // + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // CHECK: load i16, i16* [[AA_CADDR]], - // CHECK: br label {{%?}}[[TERM:.+]] + // CHECK: br label {{%?}}[[TERMINATE:.+]] // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // @@ -160,19 +169,22 @@ int foo(int n) { aa += 1; } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l261}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l276}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -188,7 +200,7 @@ int foo(int n) { // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+foo.+l261]](i[[SZ]] + // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l276]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]* @@ -219,26 +231,29 @@ int foo(int n) { // CHECK-DAG: [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]], // CHECK-DAG: [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]], // - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] - // CHECK: {{call|invoke}} void [[T4]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: {{call|invoke}} void [[T3]]_worker() + // CHECK: br label {{%?}}[[EXIT:.+]] + // + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // // Use captures. // CHECK-64-DAG: load i32, i32* [[REF_A]] @@ -249,10 +264,10 @@ int foo(int n) { // CHECK-DAG: getelementptr inbounds double, double* [[REF_CN]], i[[SZ]] %{{.+}} // CHECK-DAG: getelementptr inbounds [[TT]], [[TT]]* [[REF_D]], i32 0, i32 0 // - // CHECK: br label {{%?}}[[TERM:.+]] + // CHECK: br label {{%?}}[[TERMINATE:.+]] // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // @@ -338,19 +353,22 @@ int bar(int n){ return a; } - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+l298}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+313}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -366,7 +384,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+static.+l298]](i[[SZ]] + // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l313]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] @@ -382,36 +400,37 @@ int bar(int n){ // CHECK-DAG: [[REF_AAA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AAA]] to i8* // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]], // - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] - // CHECK: {{call|invoke}} void [[T5]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: {{call|invoke}} void [[T4]]_worker() + // CHECK: br label {{%?}}[[EXIT:.+]] // - // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // + // CHECK: [[MASTER]] + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // CHECK-64-DAG: load i32, i32* [[REF_A]] // CHECK-32-DAG: load i32, i32* [[LOCAL_A]] // CHECK-DAG: load i16, i16* [[REF_AA]] // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2 + // CHECK: br label {{%?}}[[TERMINATE:.+]] // - // CHECK: br label {{%?}}[[TERM:.+]] - // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // @@ -420,19 +439,22 @@ int bar(int n){ - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l316}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l331}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -448,7 +470,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+S1.+l316]]( + // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l331]]( // Create local storage for each capture. // CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]* // CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]] @@ -466,35 +488,39 @@ int bar(int n){ // CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA1]], // CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[LOCAL_VLA2]], // CHECK-DAG: [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]], - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] - // CHECK: {{call|invoke}} void [[T6]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: {{call|invoke}} void [[T5]]_worker() + // CHECK: br label {{%?}}[[EXIT:.+]] + // + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // Use captures. // CHECK-DAG: getelementptr inbounds [[S1]], [[S1]]* [[REF_THIS]], i32 0, i32 0 // CHECK-64-DAG:load i32, i32* [[REF_B]] // CHECK-32-DAG:load i32, i32* [[LOCAL_B]] // CHECK-DAG: getelementptr inbounds i16, i16* [[REF_C]], i[[SZ]] %{{.+}} - // CHECK: br label {{%?}}[[TERM:.+]] + // CHECK: br label {{%?}}[[TERMINATE:.+]] // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // @@ -503,19 +529,22 @@ int bar(int n){ - // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l281}}_worker() + // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l296}}_worker() + // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, + // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, + // CHECK: store i8* null, i8** [[OMP_WORK_FN]], + // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]], // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: [[WORK:%.+]] = load i64, i64 addrspace(3)* [[OMP_WID]], - // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i64 [[WORK]], 0 + // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], + // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] // // CHECK: [[SEL_WORKERS]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[NT:%.+]] = load i32, i32 addrspace(3)* [[OMP_NT]] - // CHECK: [[IS_ACTIVE:%.+]] = icmp slt i32 [[TID]], [[NT]] + // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]], + // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0 // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]] // // CHECK: [[EXEC_PARALLEL]] @@ -531,7 +560,7 @@ int bar(int n){ // CHECK: [[EXIT]] // CHECK: ret void - // CHECK: define {{.*}}void [[T7:@__omp_offloading_.+template.+l281]](i[[SZ]] + // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l296]](i[[SZ]] // Create local storage for each capture. // CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]] // CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]] @@ -544,36 +573,39 @@ int bar(int n){ // CHECK-DAG: [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16* // CHECK-DAG: [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]], // - // CHECK: [[NTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() - // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() - // CHECK: [[A:%.+]] = sub i32 [[WS]], 1 - // CHECK: [[B:%.+]] = sub i32 [[NTID]], 1 - // CHECK: [[MID:%.+]] = and i32 [[B]], - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: [[EXCESS:%.+]] = icmp ugt i32 [[TID]], [[MID]] - // CHECK: br i1 [[EXCESS]], label {{%?}}[[EXIT:.+]], label {{%?}}[[CHECK_WORKER:.+]] - // - // CHECK: [[CHECK_WORKER]] - // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[MID]] - // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[MASTER:.+]] + // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]] + // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]] + // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]] // // CHECK: [[WORKER]] - // CHECK: {{call|invoke}} void [[T7]]_worker() - // CHECK: br label {{%?}}[[EXIT]] + // CHECK: {{call|invoke}} void [[T6]]_worker() + // CHECK: br label {{%?}}[[EXIT:.+]] + // + // CHECK: [[CHECK_MASTER]] + // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], + // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]] // // CHECK: [[MASTER]] - // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() - // CHECK: call void @__kmpc_kernel_init(i32 0, i32 [[TID]]) + // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() + // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]] + // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // // CHECK-64-DAG: load i32, i32* [[REF_A]] // CHECK-32-DAG: load i32, i32* [[LOCAL_A]] // CHECK-DAG: load i16, i16* [[REF_AA]] // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2 // - // CHECK: br label {{%?}}[[TERM:.+]] + // CHECK: br label {{%?}}[[TERMINATE:.+]] // - // CHECK: [[TERM]] - // CHECK: store i64 0, i64 addrspace(3)* [[OMP_WID]], + // CHECK: [[TERMINATE]] + // CHECK: call void @__kmpc_kernel_deinit() // CHECK: call void @llvm.nvvm.barrier0() // CHECK: br label {{%?}}[[EXIT]] // diff --git a/test/OpenMP/target_codegen.cpp b/test/OpenMP/target_codegen.cpp index f263ebdd2fe3..b5e4b07cce04 100644 --- a/test/OpenMP/target_codegen.cpp +++ b/test/OpenMP/target_codegen.cpp @@ -22,11 +22,11 @@ // CHECK-DAG: [[TT:%.+]] = type { i64, i8 } // CHECK-DAG: [[S1:%.+]] = type { double } -// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] } +// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } -// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}} } +// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } // We have 8 target regions, but only 7 that actually will generate offloading // code, only 6 will have mapped arguments, and only 4 have all-constant map diff --git a/test/OpenMP/target_codegen_registration.cpp b/test/OpenMP/target_codegen_registration.cpp index a440faff9158..f2721b77fec0 100644 --- a/test/OpenMP/target_codegen_registration.cpp +++ b/test/OpenMP/target_codegen_registration.cpp @@ -30,11 +30,11 @@ // CHECK-DAG: [[SE:%.+]] = type { [64 x i32] } // CHECK-DAG: [[ST1:%.+]] = type { [228 x i32] } // CHECK-DAG: [[ST2:%.+]] = type { [1128 x i32] } -// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] } +// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } -// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] } +// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } // CHECK-DAG: [[A1:@.+]] = internal global [[SA]] // CHECK-DAG: [[A2:@.+]] = global [[SA]] @@ -100,54 +100,54 @@ // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00" -// CHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00" -// CHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00" -// CHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00" -// CHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00" -// CHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00" -// CHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00" -// CHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00" -// CHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00" -// CHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00" -// CHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00" -// CHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00" -// CHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// CHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00" -// TCHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00" -// TCHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00" -// TCHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00" -// TCHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00" -// TCHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00" -// TCHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00" -// TCHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00" -// TCHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00" -// TCHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00" -// TCHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00" -// TCHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // TCHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00" -// TCHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1 +// TCHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] diff --git a/test/OpenMP/teams_distribute_collapse_messages.cpp b/test/OpenMP/teams_distribute_collapse_messages.cpp index 9ce58e0b0650..37c10e5986bf 100644 --- a/test/OpenMP/teams_distribute_collapse_messages.cpp +++ b/test/OpenMP/teams_distribute_collapse_messages.cpp @@ -66,7 +66,8 @@ T tmain(T argc, S **argv) { //expected-note 2 {{declared here}} for (int i = ST; i < N; i++) argv[0][i] = argv[0][i] - argv[0][i-ST]; -#pragma omp distribute collapse (S) // expected-error {{'S' does not refer to a value}} +#pragma omp target +#pragma omp teams distribute collapse (S) // expected-error {{'S' does not refer to a value}} for (int i = ST; i < N; i++) argv[0][i] = argv[0][i] - argv[0][i-ST]; diff --git a/test/Preprocessor/cuda-types.cu b/test/Preprocessor/cuda-types.cu index 2b6160b8d6c7..5f7b91655cdf 100644 --- a/test/Preprocessor/cuda-types.cu +++ b/test/Preprocessor/cuda-types.cu @@ -28,3 +28,19 @@ // RUN: | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \ // RUN: | grep -v '__LDBL\|_LONG_DOUBLE' > %T/powerpc64-device-defines-filtered // RUN: diff %T/powerpc64-host-defines-filtered %T/powerpc64-device-defines-filtered + +// RUN: %clang --cuda-host-only -nocudainc -target i386-windows-msvc -x cuda -E -dM -o - /dev/null \ +// RUN: | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \ +// RUN: | grep -v '__LDBL\|_LONG_DOUBLE' > %T/i386-msvc-host-defines-filtered +// RUN: %clang --cuda-device-only -nocudainc -nocudalib -target i386-windows-msvc -x cuda -E -dM -o - /dev/null \ +// RUN: | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \ +// RUN: | grep -v '__LDBL\|_LONG_DOUBLE' > %T/i386-msvc-device-defines-filtered +// RUN: diff %T/i386-msvc-host-defines-filtered %T/i386-msvc-device-defines-filtered + +// RUN: %clang --cuda-host-only -nocudainc -target x86_64-windows-msvc -x cuda -E -dM -o - /dev/null \ +// RUN: | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \ +// RUN: | grep -v '__LDBL\|_LONG_DOUBLE' > %T/x86_64-msvc-host-defines-filtered +// RUN: %clang --cuda-device-only -nocudainc -nocudalib -target x86_64-windows-msvc -x cuda -E -dM -o - /dev/null \ +// RUN: | grep 'define __[^ ]*\(TYPE\|MAX\|SIZEOF|WIDTH\)\|define __GCC_ATOMIC' \ +// RUN: | grep -v '__LDBL\|_LONG_DOUBLE' > %T/x86_64-msvc-device-defines-filtered +// RUN: diff %T/x86_64-msvc-host-defines-filtered %T/x86_64-msvc-device-defines-filtered diff --git a/test/Preprocessor/init.c b/test/Preprocessor/init.c index b003404df6ff..8b8901931e7a 100644 --- a/test/Preprocessor/init.c +++ b/test/Preprocessor/init.c @@ -9189,3 +9189,174 @@ // RUN: %clang_cc1 -E -dM -ffreestanding -triple x86_64-windows-cygnus < /dev/null | FileCheck -match-full-lines -check-prefix CYGWIN-X64 %s // CYGWIN-X64: #define __USER_LABEL_PREFIX__ +// RUN: %clang_cc1 -E -dM -ffreestanding -triple=avr \ +// RUN: < /dev/null \ +// RUN: | FileCheck -match-full-lines -check-prefix=AVR %s +// +// AVR:#define __ATOMIC_ACQUIRE 2 +// AVR:#define __ATOMIC_ACQ_REL 4 +// AVR:#define __ATOMIC_CONSUME 1 +// AVR:#define __ATOMIC_RELAXED 0 +// AVR:#define __ATOMIC_RELEASE 3 +// AVR:#define __ATOMIC_SEQ_CST 5 +// AVR:#define __AVR__ 1 +// AVR:#define __BIGGEST_ALIGNMENT__ 1 +// AVR:#define __BYTE_ORDER__ __ORDER_LITTLE_ENDIAN__ +// AVR:#define __CHAR16_TYPE__ unsigned int +// AVR:#define __CHAR32_TYPE__ long unsigned int +// AVR:#define __CHAR_BIT__ 8 +// AVR:#define __DBL_DECIMAL_DIG__ 9 +// AVR:#define __DBL_DENORM_MIN__ 1.40129846e-45 +// AVR:#define __DBL_DIG__ 6 +// AVR:#define __DBL_EPSILON__ 1.19209290e-7 +// AVR:#define __DBL_HAS_DENORM__ 1 +// AVR:#define __DBL_HAS_INFINITY__ 1 +// AVR:#define __DBL_HAS_QUIET_NAN__ 1 +// AVR:#define __DBL_MANT_DIG__ 24 +// AVR:#define __DBL_MAX_10_EXP__ 38 +// AVR:#define __DBL_MAX_EXP__ 128 +// AVR:#define __DBL_MAX__ 3.40282347e+38 +// AVR:#define __DBL_MIN_10_EXP__ (-37) +// AVR:#define __DBL_MIN_EXP__ (-125) +// AVR:#define __DBL_MIN__ 1.17549435e-38 +// AVR:#define __FINITE_MATH_ONLY__ 0 +// AVR:#define __FLT_DECIMAL_DIG__ 9 +// AVR:#define __FLT_DENORM_MIN__ 1.40129846e-45F +// AVR:#define __FLT_DIG__ 6 +// AVR:#define __FLT_EPSILON__ 1.19209290e-7F +// AVR:#define __FLT_EVAL_METHOD__ 0 +// AVR:#define __FLT_HAS_DENORM__ 1 +// AVR:#define __FLT_HAS_INFINITY__ 1 +// AVR:#define __FLT_HAS_QUIET_NAN__ 1 +// AVR:#define __FLT_MANT_DIG__ 24 +// AVR:#define __FLT_MAX_10_EXP__ 38 +// AVR:#define __FLT_MAX_EXP__ 128 +// AVR:#define __FLT_MAX__ 3.40282347e+38F +// AVR:#define __FLT_MIN_10_EXP__ (-37) +// AVR:#define __FLT_MIN_EXP__ (-125) +// AVR:#define __FLT_MIN__ 1.17549435e-38F +// AVR:#define __FLT_RADIX__ 2 +// AVR:#define __GCC_ATOMIC_BOOL_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_CHAR16_T_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_CHAR32_T_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_CHAR_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_INT_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_LLONG_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_LONG_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_POINTER_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_SHORT_LOCK_FREE 1 +// AVR:#define __GCC_ATOMIC_TEST_AND_SET_TRUEVAL 1 +// AVR:#define __GCC_ATOMIC_WCHAR_T_LOCK_FREE 1 +// AVR:#define __GXX_ABI_VERSION 1002 +// AVR:#define __INT16_C_SUFFIX__ +// AVR:#define __INT16_MAX__ 32767 +// AVR:#define __INT16_TYPE__ short +// AVR:#define __INT32_C_SUFFIX__ L +// AVR:#define __INT32_MAX__ 2147483647L +// AVR:#define __INT32_TYPE__ long int +// AVR:#define __INT64_C_SUFFIX__ LL +// AVR:#define __INT64_MAX__ 9223372036854775807LL +// AVR:#define __INT64_TYPE__ long long int +// AVR:#define __INT8_C_SUFFIX__ +// AVR:#define __INT8_MAX__ 127 +// AVR:#define __INT8_TYPE__ signed char +// AVR:#define __INTMAX_C_SUFFIX__ LL +// AVR:#define __INTMAX_MAX__ 9223372036854775807LL +// AVR:#define __INTMAX_TYPE__ long long int +// AVR:#define __INTPTR_MAX__ 32767 +// AVR:#define __INTPTR_TYPE__ int +// AVR:#define __INT_FAST16_MAX__ 32767 +// AVR:#define __INT_FAST16_TYPE__ int +// AVR:#define __INT_FAST32_MAX__ 2147483647L +// AVR:#define __INT_FAST32_TYPE__ long int +// AVR:#define __INT_FAST64_MAX__ 9223372036854775807LL +// AVR:#define __INT_FAST64_TYPE__ long long int +// AVR:#define __INT_FAST8_MAX__ 127 +// AVR:#define __INT_FAST8_TYPE__ signed char +// AVR:#define __INT_LEAST16_MAX__ 32767 +// AVR:#define __INT_LEAST16_TYPE__ int +// AVR:#define __INT_LEAST32_MAX__ 2147483647L +// AVR:#define __INT_LEAST32_TYPE__ long int +// AVR:#define __INT_LEAST64_MAX__ 9223372036854775807LL +// AVR:#define __INT_LEAST64_TYPE__ long long int +// AVR:#define __INT_LEAST8_MAX__ 127 +// AVR:#define __INT_LEAST8_TYPE__ signed char +// AVR:#define __INT_MAX__ 32767 +// AVR:#define __LDBL_DECIMAL_DIG__ 9 +// AVR:#define __LDBL_DENORM_MIN__ 1.40129846e-45L +// AVR:#define __LDBL_DIG__ 6 +// AVR:#define __LDBL_EPSILON__ 1.19209290e-7L +// AVR:#define __LDBL_HAS_DENORM__ 1 +// AVR:#define __LDBL_HAS_INFINITY__ 1 +// AVR:#define __LDBL_HAS_QUIET_NAN__ 1 +// AVR:#define __LDBL_MANT_DIG__ 24 +// AVR:#define __LDBL_MAX_10_EXP__ 38 +// AVR:#define __LDBL_MAX_EXP__ 128 +// AVR:#define __LDBL_MAX__ 3.40282347e+38L +// AVR:#define __LDBL_MIN_10_EXP__ (-37) +// AVR:#define __LDBL_MIN_EXP__ (-125) +// AVR:#define __LDBL_MIN__ 1.17549435e-38L +// AVR:#define __LONG_LONG_MAX__ 9223372036854775807LL +// AVR:#define __LONG_MAX__ 2147483647L +// AVR:#define __NO_INLINE__ 1 +// AVR:#define __ORDER_BIG_ENDIAN__ 4321 +// AVR:#define __ORDER_LITTLE_ENDIAN__ 1234 +// AVR:#define __ORDER_PDP_ENDIAN__ 3412 +// AVR:#define __PRAGMA_REDEFINE_EXTNAME 1 +// AVR:#define __PTRDIFF_MAX__ 32767 +// AVR:#define __PTRDIFF_TYPE__ int +// AVR:#define __SCHAR_MAX__ 127 +// AVR:#define __SHRT_MAX__ 32767 +// AVR:#define __SIG_ATOMIC_MAX__ 127 +// AVR:#define __SIG_ATOMIC_WIDTH__ 8 +// AVR:#define __SIZEOF_DOUBLE__ 4 +// AVR:#define __SIZEOF_FLOAT__ 4 +// AVR:#define __SIZEOF_INT__ 2 +// AVR:#define __SIZEOF_LONG_DOUBLE__ 4 +// AVR:#define __SIZEOF_LONG_LONG__ 8 +// AVR:#define __SIZEOF_LONG__ 4 +// AVR:#define __SIZEOF_POINTER__ 2 +// AVR:#define __SIZEOF_PTRDIFF_T__ 2 +// AVR:#define __SIZEOF_SHORT__ 2 +// AVR:#define __SIZEOF_SIZE_T__ 2 +// AVR:#define __SIZEOF_WCHAR_T__ 2 +// AVR:#define __SIZEOF_WINT_T__ 2 +// AVR:#define __SIZE_MAX__ 65535U +// AVR:#define __SIZE_TYPE__ unsigned int +// AVR:#define __STDC__ 1 +// AVR:#define __UINT16_MAX__ 65535U +// AVR:#define __UINT16_TYPE__ unsigned short +// AVR:#define __UINT32_C_SUFFIX__ UL +// AVR:#define __UINT32_MAX__ 4294967295UL +// AVR:#define __UINT32_TYPE__ long unsigned int +// AVR:#define __UINT64_C_SUFFIX__ ULL +// AVR:#define __UINT64_MAX__ 18446744073709551615ULL +// AVR:#define __UINT64_TYPE__ long long unsigned int +// AVR:#define __UINT8_C_SUFFIX__ +// AVR:#define __UINT8_MAX__ 255 +// AVR:#define __UINT8_TYPE__ unsigned char +// AVR:#define __UINTMAX_C_SUFFIX__ ULL +// AVR:#define __UINTMAX_MAX__ 18446744073709551615ULL +// AVR:#define __UINTMAX_TYPE__ long long unsigned int +// AVR:#define __UINTPTR_MAX__ 65535U +// AVR:#define __UINTPTR_TYPE__ unsigned int +// AVR:#define __UINT_FAST16_MAX__ 65535U +// AVR:#define __UINT_FAST16_TYPE__ unsigned int +// AVR:#define __UINT_FAST32_MAX__ 4294967295UL +// AVR:#define __UINT_FAST32_TYPE__ long unsigned int +// AVR:#define __UINT_FAST64_MAX__ 18446744073709551615ULL +// AVR:#define __UINT_FAST64_TYPE__ long long unsigned int +// AVR:#define __UINT_FAST8_MAX__ 255 +// AVR:#define __UINT_FAST8_TYPE__ unsigned char +// AVR:#define __UINT_LEAST16_MAX__ 65535U +// AVR:#define __UINT_LEAST16_TYPE__ unsigned int +// AVR:#define __UINT_LEAST32_MAX__ 4294967295UL +// AVR:#define __UINT_LEAST32_TYPE__ long unsigned int +// AVR:#define __UINT_LEAST64_MAX__ 18446744073709551615ULL +// AVR:#define __UINT_LEAST64_TYPE__ long long unsigned int +// AVR:#define __UINT_LEAST8_MAX__ 255 +// AVR:#define __UINT_LEAST8_TYPE__ unsigned char +// AVR:#define __USER_LABEL_PREFIX__ +// AVR:#define __WCHAR_MAX__ 32767 +// AVR:#define __WCHAR_TYPE__ int +// AVR:#define __WINT_TYPE__ int diff --git a/test/Sema/warn-cast-align.c b/test/Sema/warn-cast-align.c index e8f85bc14d8d..389c0c17d2f7 100644 --- a/test/Sema/warn-cast-align.c +++ b/test/Sema/warn-cast-align.c @@ -59,3 +59,11 @@ void test4() { i = (int *)&s.s0; i = (int *)a; } + +// No warnings. +typedef int (*FnTy)(void); +unsigned int func5(void); + +FnTy test5(void) { + return (FnTy)&func5; +} diff --git a/test/Sema/warn-strict-prototypes.m b/test/Sema/warn-strict-prototypes.m index cbb01a1f7b21..4567dab01930 100644 --- a/test/Sema/warn-strict-prototypes.m +++ b/test/Sema/warn-strict-prototypes.m @@ -14,7 +14,8 @@ void foo() { void (^block)() = // expected-warning {{this function declaration is not a prototype}} ^void(int arg) { // no warning }; - void (^block2)(void) = // no warning - ^void() { // expected-warning {{this function declaration is not a prototype}} + void (^block2)(void) = ^void() { // no warning + }; + void (^block3)(void) = ^ { // no warning }; } diff --git a/test/Sema/warn-thread-safety-analysis.c b/test/Sema/warn-thread-safety-analysis.c index a0c4026b9136..425ce4c196a6 100644 --- a/test/Sema/warn-thread-safety-analysis.c +++ b/test/Sema/warn-thread-safety-analysis.c @@ -127,3 +127,7 @@ int main() { return 0; } + +// We had a problem where we'd skip all attributes that follow a late-parsed +// attribute in a single __attribute__. +void run() __attribute__((guarded_by(mu1), guarded_by(mu1))); // expected-warning 2{{only applies to fields and global variables}} diff --git a/test/SemaCUDA/attr-declspec.cu b/test/SemaCUDA/attr-declspec.cu new file mode 100644 index 000000000000..dda12ce8a51f --- /dev/null +++ b/test/SemaCUDA/attr-declspec.cu @@ -0,0 +1,34 @@ +// Test the __declspec spellings of CUDA attributes. +// +// RUN: %clang_cc1 -fsyntax-only -fms-extensions -verify %s +// RUN: %clang_cc1 -fsyntax-only -fms-extensions -fcuda-is-device -verify %s +// Now pretend that we're compiling a C file. There should be warnings. +// RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s + +#if defined(EXPECT_WARNINGS) +// expected-warning@+12 {{'__device__' attribute ignored}} +// expected-warning@+12 {{'__global__' attribute ignored}} +// expected-warning@+12 {{'__constant__' attribute ignored}} +// expected-warning@+12 {{'__shared__' attribute ignored}} +// expected-warning@+12 {{'__host__' attribute ignored}} +// +// (Currently we don't for the other attributes. They are implemented with +// IgnoredAttr, which is ignored irrespective of any LangOpts.) +#else +// expected-no-diagnostics +#endif + +__declspec(__device__) void f_device(); +__declspec(__global__) void f_global(); +__declspec(__constant__) int* g_constant; +__declspec(__shared__) float *g_shared; +__declspec(__host__) void f_host(); +__declspec(__device_builtin__) void f_device_builtin(); +typedef __declspec(__device_builtin__) const void *t_device_builtin; +enum __declspec(__device_builtin__) e_device_builtin {E}; +__declspec(__device_builtin__) int v_device_builtin; +__declspec(__cudart_builtin__) void f_cudart_builtin(); +__declspec(__device_builtin_surface_type__) unsigned long long surface_var; +__declspec(__device_builtin_texture_type__) unsigned long long texture_var; + +// Note that there's no __declspec spelling of nv_weak. diff --git a/test/SemaCUDA/cuda-inherits-calling-conv.cu b/test/SemaCUDA/cuda-inherits-calling-conv.cu new file mode 100644 index 000000000000..67c438fa621b --- /dev/null +++ b/test/SemaCUDA/cuda-inherits-calling-conv.cu @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -std=c++11 -triple i386-windows-msvc \ +// RUN: -aux-triple nvptx-nvidia-cuda -fsyntax-only -verify %s + +// RUN: %clang_cc1 -std=c++11 -triple nvptx-nvidia-cuda \ +// RUN: -aux-triple i386-windows-msvc -fsyntax-only \ +// RUN: -fcuda-is-device -verify %s + +// RUN: %clang_cc1 -std=c++11 -triple nvptx-nvidia-cuda \ +// RUN: -aux-triple x86_64-linux-gnu -fsyntax-only \ +// RUN: -fcuda-is-device -verify -verify-ignore-unexpected=note \ +// RUN: -DEXPECT_ERR %s + +// CUDA device code should inherit the host's calling conventions. + +template <class T> +struct Foo; + +template <class T> +struct Foo<T()> {}; + +// On x86_64-linux-gnu, this is a redefinition of the template, because the +// __fastcall calling convention doesn't exist (and is therefore ignored). +#ifndef EXPECT_ERR +// expected-no-diagnostics +#else +// expected-error@+4 {{redefinition of 'Foo}} +// expected-warning@+3 {{calling convention '__fastcall' ignored}} +#endif +template <class T> +struct Foo<T __fastcall()> {}; diff --git a/test/SemaCXX/constant-expression-cxx11.cpp b/test/SemaCXX/constant-expression-cxx11.cpp index 581a524339e7..884f2f30c42f 100644 --- a/test/SemaCXX/constant-expression-cxx11.cpp +++ b/test/SemaCXX/constant-expression-cxx11.cpp @@ -1725,7 +1725,7 @@ namespace AfterError { constexpr int error() { // expected-error {{no return statement}} return foobar; // expected-error {{undeclared identifier}} } - constexpr int k = error(); // expected-error {{must be initialized by a constant expression}} + constexpr int k = error(); } namespace std { @@ -2030,7 +2030,7 @@ namespace PR21786 { namespace PR21859 { constexpr int Fun() { return; } // expected-error {{non-void constexpr function 'Fun' should return a value}} - constexpr int Var = Fun(); // expected-error {{constexpr variable 'Var' must be initialized by a constant expression}} + constexpr int Var = Fun(); } struct InvalidRedef { diff --git a/test/SemaCXX/conversion-function.cpp b/test/SemaCXX/conversion-function.cpp index c725a0d5b7c1..531de818b680 100644 --- a/test/SemaCXX/conversion-function.cpp +++ b/test/SemaCXX/conversion-function.cpp @@ -440,7 +440,7 @@ namespace PR18234 { #endif } a; A::S s = a; // expected-error {{no viable conversion from 'struct A' to 'A::S'}} - A::E e = a; // expected-note {{here}} + A::E e = a; bool k1 = e == A::e; // expected-error {{no member named 'e'}} bool k2 = e.n == 0; } diff --git a/test/SemaCXX/cxx0x-initializer-stdinitializerlist.cpp b/test/SemaCXX/cxx0x-initializer-stdinitializerlist.cpp index 75c6734bce30..9b8fadd2f522 100644 --- a/test/SemaCXX/cxx0x-initializer-stdinitializerlist.cpp +++ b/test/SemaCXX/cxx0x-initializer-stdinitializerlist.cpp @@ -105,6 +105,7 @@ T deduce_ref(const std::initializer_list<T>&); // expected-note {{conflicting ty template<typename T, typename U> struct pair { pair(...); }; template<typename T> void deduce_pairs(std::initializer_list<pair<T, typename T::type>>); +// expected-note@-1 {{deduced type 'pair<[...], typename WithIntType::type>' of element of 1st parameter does not match adjusted type 'pair<[...], float>' of element of argument [with T = WithIntType]}} struct WithIntType { typedef int type; }; template<typename ...T> void deduce_after_init_list_in_pack(void (*)(T...), T...); // expected-note {{<int, int> vs. <(no value), double>}} @@ -123,7 +124,7 @@ void argument_deduction() { pair<WithIntType, int> pi; pair<WithIntType, float> pf; deduce_pairs({pi, pi, pi}); // ok - deduce_pairs({pi, pf, pi}); // FIXME: This should be rejected, as we fail to produce a type that exactly matches the argument type. + deduce_pairs({pi, pf, pi}); // expected-error {{no matching function}} deduce_after_init_list_in_pack((void(*)(int,int))0, {}, 0); deduce_after_init_list_in_pack((void(*)(int,int))0, {}, 0.0); // expected-error {{no matching function}} @@ -298,9 +299,18 @@ namespace TemporaryInitListSourceRange_PR22367 { namespace ParameterPackNestedInitializerLists_PR23904c3 { template <typename ...T> - void f(std::initializer_list<std::initializer_list<T>> ...tt); + void f(std::initializer_list<std::initializer_list<T>> ...tt); // expected-note 2{{conflicting}} expected-note {{incomplete pack}} - void foo() { f({{0}}, {{'\0'}}); } + void foo() { + f({{0}}, {{'\0'}}); // ok, T = <int, char> + f({{0}, {'\0'}}); // expected-error {{no match}} + f({{0, '\0'}}); // expected-error {{no match}} + + f({{0}}, {{{}}}); // expected-error {{no match}} + f({{0}}, {{{}, '\0'}}); // ok, T = <int, char> + f({{0}, {{}}}); // ok, T = <int> + f({{0, {}}}); // ok, T = <int> + } } namespace update_rbrace_loc_crash { @@ -327,3 +337,13 @@ namespace update_rbrace_loc_crash { Explode<ContainsIncomplete, 4>([](int) {}); } } + +namespace no_conversion_after_auto_list_deduction { + // We used to deduce 'auto' == 'std::initializer_list<X>' here, and then + // incorrectly accept the declaration of 'x'. + struct X { using T = std::initializer_list<X> X::*; operator T(); }; + auto X::*x = { X() }; // expected-error {{from initializer list}} + + struct Y { using T = std::initializer_list<Y>(*)(); operator T(); }; + auto (*y)() = { Y() }; // expected-error {{from initializer list}} +} diff --git a/test/SemaCXX/cxx1z-decomposition.cpp b/test/SemaCXX/cxx1z-decomposition.cpp index 735a9e1dfee0..d457ace5d844 100644 --- a/test/SemaCXX/cxx1z-decomposition.cpp +++ b/test/SemaCXX/cxx1z-decomposition.cpp @@ -65,4 +65,9 @@ void for_range() { } } +int error_recovery() { + auto [foobar]; // expected-error {{requires an initializer}} + return foobar_; // expected-error {{undeclared identifier 'foobar_'}} +} + // FIXME: by-value array copies diff --git a/test/SemaCXX/default-arg-closures.cpp b/test/SemaCXX/default-arg-closures.cpp index e076cc05cd20..676bd486105f 100644 --- a/test/SemaCXX/default-arg-closures.cpp +++ b/test/SemaCXX/default-arg-closures.cpp @@ -4,16 +4,15 @@ // instantiating and checking the semantics of default arguments. Make sure we // do that right. -// FIXME: Don't diagnose this issue twice. template <typename T> -struct DependentDefaultCtorArg { // expected-note {{in instantiation of default function argument}} - // expected-error@+1 2 {{type 'int' cannot be used prior to '::' because it has no members}} +struct DependentDefaultCtorArg { + // expected-error@+1 {{type 'int' cannot be used prior to '::' because it has no members}} DependentDefaultCtorArg(int n = T::error); }; struct __declspec(dllexport) // expected-note {{due to 'ExportDefaultCtorClosure' being dllexported}} -ExportDefaultCtorClosure // expected-note {{implicit default constructor for 'ExportDefaultCtorClosure' first required here}} -: DependentDefaultCtorArg<int> // expected-note {{in instantiation of template class}} +ExportDefaultCtorClosure // expected-note {{in instantiation of default function argument expression for 'DependentDefaultCtorArg<int>' required here}} expected-note {{implicit default constructor for 'ExportDefaultCtorClosure' first required here}} +: DependentDefaultCtorArg<int> {}; template <typename T> diff --git a/test/SemaCXX/dllexport.cpp b/test/SemaCXX/dllexport.cpp index b4850fc03d9b..a3fed70ec958 100644 --- a/test/SemaCXX/dllexport.cpp +++ b/test/SemaCXX/dllexport.cpp @@ -741,6 +741,27 @@ struct __declspec(dllexport) ClassWithMultipleDefaultCtors { ClassWithMultipleDefaultCtors(int = 40) {} // expected-error{{'__declspec(dllexport)' cannot be applied to more than one default constructor}} ClassWithMultipleDefaultCtors(int = 30, ...) {} // expected-note{{declared here}} }; +template <typename T> +struct ClassTemplateWithMultipleDefaultCtors { + __declspec(dllexport) ClassTemplateWithMultipleDefaultCtors(int = 40) {} // expected-error{{'__declspec(dllexport)' cannot be applied to more than one default constructor}} + __declspec(dllexport) ClassTemplateWithMultipleDefaultCtors(int = 30, ...) {} // expected-note{{declared here}} +}; + +template <typename T> struct HasDefaults { + HasDefaults(int x = sizeof(T)) {} // expected-error {{invalid application of 'sizeof'}} +}; +template struct __declspec(dllexport) HasDefaults<char>; + +template struct +__declspec(dllexport) // expected-note {{in instantiation of default function argument expression for 'HasDefaults<void>' required here}} +HasDefaults<void>; // expected-note {{in instantiation of member function 'HasDefaults<void>::HasDefaults' requested here}} + +template <typename T> struct HasDefaults2 { + __declspec(dllexport) // expected-note {{in instantiation of default function argument expression for 'HasDefaults2<void>' required here}} + HasDefaults2(int x = sizeof(T)) {} // expected-error {{invalid application of 'sizeof'}} +}; +template struct HasDefaults2<void>; // expected-note {{in instantiation of member function 'HasDefaults2<void>::HasDefaults2' requested here}} + #endif //===----------------------------------------------------------------------===// diff --git a/test/SemaCXX/type-definition-in-specifier.cpp b/test/SemaCXX/type-definition-in-specifier.cpp index 74ba058b4f12..2da649fdb0b8 100644 --- a/test/SemaCXX/type-definition-in-specifier.cpp +++ b/test/SemaCXX/type-definition-in-specifier.cpp @@ -59,10 +59,8 @@ struct s19018b { }; struct pr18963 { - short bar5 (struct foo4 {} bar2); // expected-error{{'foo4' cannot be defined in a parameter type}} \ - // expected-note{{declared here}} - - long foo5 (float foo6 = foo4); // expected-error{{'foo4' does not refer to a value}} + short bar5 (struct foo4 {} bar2); // expected-error{{'foo4' cannot be defined in a parameter type}} + long foo5 (float foo6 = foo4); }; // expected-error@+2 {{cannot be defined in a parameter type}} diff --git a/test/SemaObjC/block-omitted-return-type.m b/test/SemaObjC/block-omitted-return-type.m index 20e32e01865e..93d5e05ea282 100644 --- a/test/SemaObjC/block-omitted-return-type.m +++ b/test/SemaObjC/block-omitted-return-type.m @@ -24,7 +24,7 @@ return; }; void (^simpleBlock5)() = ^ const void { //expected-error {{incompatible block pointer types initializing 'void (^)()' with an expression of type 'const void (^)(void)'}} - return; + return; // expected-warning@-1 {{function cannot return qualified void type 'const void'}} }; void (^simpleBlock6)() = ^ const (void) { //expected-warning {{'const' qualifier on omitted return type '<dependent type>' has no effect}} return; diff --git a/test/SemaOpenCL/extensions.cl b/test/SemaOpenCL/extensions.cl index c27f3397cd79..6afb11e42a6a 100644 --- a/test/SemaOpenCL/extensions.cl +++ b/test/SemaOpenCL/extensions.cl @@ -22,6 +22,17 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -verify -pedantic -fsyntax-only -cl-ext=-all -cl-ext=+cl_khr_fp64 -cl-ext=+cl_khr_fp16 -cl-ext=-cl_khr_fp64 -DNOFP64 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -verify -pedantic -fsyntax-only -cl-ext=-all -cl-ext=+cl_khr_fp64,-cl_khr_fp64,+cl_khr_fp16 -DNOFP64 +// Test with -finclude-default-header, which includes opencl-c.h. opencl-c.h +// disables all extensions by default, but supported core extensions for a +// particular OpenCL version must be re-enabled (for example, cl_khr_fp64 is +// enabled by default with -cl-std=CL2.0). +// +// RUN: %clang_cc1 %s -triple amdgcn-unknown-unknown -verify -pedantic -fsyntax-only -cl-std=CL2.0 -finclude-default-header + +#ifdef _OPENCL_H_ +// expected-no-diagnostics +#endif + #ifdef FP64 // expected-no-diagnostics #endif @@ -33,6 +44,7 @@ void f1(double da) { // expected-error {{type 'double' requires cl_khr_fp64 exte } #endif +#ifndef _OPENCL_H_ int isnan(float x) { return __builtin_isnan(x); } @@ -40,6 +52,7 @@ int isnan(float x) { int isfinite(float x) { return __builtin_isfinite(x); } +#endif #pragma OPENCL EXTENSION cl_khr_fp64 : enable #ifdef NOFP64 diff --git a/test/SemaTemplate/deduction.cpp b/test/SemaTemplate/deduction.cpp index 5695cab9a27e..2275a8b3b7ad 100644 --- a/test/SemaTemplate/deduction.cpp +++ b/test/SemaTemplate/deduction.cpp @@ -407,3 +407,38 @@ namespace overload_vs_pack { void test() { j(x, f, x); } } } + +namespace b29946541 { + template<typename> class A {}; + template<typename T, typename U, template<typename, typename> class C> + void f(C<T, U>); // expected-note {{failed template argument deduction}} + void g(A<int> a) { f(a); } // expected-error {{no match}} +} + +namespace deduction_from_empty_list { + template<int M, int N = 5> void f(int (&&)[N], int (&&)[N]) { // expected-note {{1 vs. 2}} + static_assert(M == N, ""); + } + + void test() { + f<5>({}, {}); + f<1>({}, {0}); + f<1>({0}, {}); + f<1>({0}, {0}); + f<1>({0}, {0, 1}); // expected-error {{no matching}} + } +} + +namespace check_extended_pack { + template<typename T> struct X { typedef int type; }; + template<typename ...T> void f(typename X<T>::type...); + template<typename T> void f(T, int, int); + void g() { + f<int>(0, 0, 0); + } + + template<int, int*> struct Y {}; + template<int ...N> void g(Y<N...>); // expected-note {{deduced non-type template argument does not have the same type as the corresponding template parameter ('int *' vs 'int')}} + int n; + void h() { g<0>(Y<0, &n>()); } // expected-error {{no matching function}} +} diff --git a/test/SemaTemplate/instantiate-local-class.cpp b/test/SemaTemplate/instantiate-local-class.cpp index a61af7a5af38..eaff4c4bbc8d 100644 --- a/test/SemaTemplate/instantiate-local-class.cpp +++ b/test/SemaTemplate/instantiate-local-class.cpp @@ -475,3 +475,14 @@ namespace rdar23721638 { } template void bar<A>(); // expected-note {{in instantiation}} } + +namespace anon_union_default_member_init { + template<typename T> void f() { + struct S { + union { + int i = 0; + }; + }; + } + void g() { f<int>(); } +} |