diff options
Diffstat (limited to 'clang/test')
73 files changed, 2905 insertions, 258 deletions
diff --git a/clang/test/APINotes/Inputs/Headers/SwiftImportAs.apinotes b/clang/test/APINotes/Inputs/Headers/SwiftImportAs.apinotes index 66fc46e..c096822 100644 --- a/clang/test/APINotes/Inputs/Headers/SwiftImportAs.apinotes +++ b/clang/test/APINotes/Inputs/Headers/SwiftImportAs.apinotes @@ -19,6 +19,10 @@ Tags: SwiftReleaseOp: release SwiftRetainOp: retain SwiftDefaultOwnership: unretained +- Name: OpaqueRefCountedType + SwiftImportAs: reference + SwiftReleaseOp: ORCRelease + SwiftRetainOp: ORCRetain - Name: NonCopyableType SwiftCopyable: false SwiftConformsTo: MySwiftModule.MySwiftNonCopyableProtocol diff --git a/clang/test/APINotes/Inputs/Headers/SwiftImportAs.h b/clang/test/APINotes/Inputs/Headers/SwiftImportAs.h index 20b8f04..5f817ac 100644 --- a/clang/test/APINotes/Inputs/Headers/SwiftImportAs.h +++ b/clang/test/APINotes/Inputs/Headers/SwiftImportAs.h @@ -23,3 +23,9 @@ struct EscapableType { int value; }; struct RefCountedTypeWithDefaultConvention {}; inline void retain(RefCountedType *x) {} inline void release(RefCountedType *x) {} + +struct OpaqueRefCountedType; +struct OpaqueRefCountedType; // redeclaration + +inline void ORCRetain(struct OpaqueRefCountedType *x); +inline void ORCRelease(struct OpaqueRefCountedType *x); diff --git a/clang/test/APINotes/swift-import-as.cpp b/clang/test/APINotes/swift-import-as.cpp index 929f924..179170f 100644 --- a/clang/test/APINotes/swift-import-as.cpp +++ b/clang/test/APINotes/swift-import-as.cpp @@ -3,6 +3,7 @@ // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter ImmortalRefType | FileCheck -check-prefix=CHECK-IMMORTAL %s // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter RefCountedType | FileCheck -check-prefix=CHECK-REF-COUNTED %s // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter RefCountedTypeWithDefaultConvention | FileCheck -check-prefix=CHECK-REF-COUNTED-DEFAULT %s +// RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter OpaqueRefCountedType | FileCheck -check-prefix=CHECK-OPAQUE-REF-COUNTED %s // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter NonCopyableType | FileCheck -check-prefix=CHECK-NON-COPYABLE %s // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter CopyableType | FileCheck -check-prefix=CHECK-COPYABLE %s // RUN: %clang_cc1 -fmodules -fblocks -fimplicit-module-maps -fmodules-cache-path=%t/ModulesCache -fdisable-module-hash -fapinotes-modules -I %S/Inputs/Headers %s -x c++ -ast-dump -ast-dump-filter NonEscapableType | FileCheck -check-prefix=CHECK-NON-ESCAPABLE %s @@ -34,6 +35,20 @@ // CHECK-REF-COUNTED-DEFAULT: SwiftAttrAttr {{.+}} <<invalid sloc>> "release:release" // CHECK-REF-COUNTED-DEFAULT: SwiftAttrAttr {{.+}} <<invalid sloc>> "returned_as_unretained_by_default" +// CHECK-OPAQUE-REF-COUNTED: Dumping OpaqueRefCountedType: +// CHECK-OPAQUE-REF-COUNTED-NEXT: CXXRecordDecl {{.+}} imported in SwiftImportAs{{.*}}struct OpaqueRefCountedType +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "import_reference" +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "retain:ORCRetain" +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "release:ORCRelease" +// CHECK-OPAQUE-REF-COUNTED-NOT: SwiftAttrAttr {{.+}} <<invalid sloc>> "release:ORCRelease" + +// CHECK-OPAQUE-REF-COUNTED: Dumping OpaqueRefCountedType: +// CHECK-OPAQUE-REF-COUNTED-NEXT: CXXRecordDecl {{.+}} imported in SwiftImportAs{{.*}}struct OpaqueRefCountedType +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "import_reference" +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "retain:ORCRetain" +// CHECK-OPAQUE-REF-COUNTED: SwiftAttrAttr {{.+}} <<invalid sloc>> "release:ORCRelease" + +// CHECK-OPAQUE-REF-COUNTED-NOT: SwiftAttrAttr {{.+}} <<invalid sloc>> "release: // CHECK-NON-COPYABLE: Dumping NonCopyableType: // CHECK-NON-COPYABLE-NEXT: CXXRecordDecl {{.+}} imported in SwiftImportAs {{.+}} struct NonCopyableType // CHECK-NON-COPYABLE: SwiftAttrAttr {{.+}} <<invalid sloc>> "conforms_to:MySwiftModule.MySwiftNonCopyableProtocol" diff --git a/clang/test/AST/ByteCode/builtin-bit-cast.cpp b/clang/test/AST/ByteCode/builtin-bit-cast.cpp index 3c5e89d7..bc356b0 100644 --- a/clang/test/AST/ByteCode/builtin-bit-cast.cpp +++ b/clang/test/AST/ByteCode/builtin-bit-cast.cpp @@ -22,6 +22,10 @@ typedef __INTPTR_TYPE__ intptr_t; static_assert(sizeof(int) == 4); static_assert(sizeof(long long) == 8); + +constexpr bool test_bad_bool = __builtin_bit_cast(bool, (char)0xff); // both-error {{must be initialized by a constant expression}} \ + // both-note {{value 255 cannot be represented in type 'bool'}} + template <class To, class From> constexpr To bit_cast(const From &from) { static_assert(sizeof(To) == sizeof(From)); diff --git a/clang/test/AST/ByteCode/unions.cpp b/clang/test/AST/ByteCode/unions.cpp index 7cfd0d6..139e318 100644 --- a/clang/test/AST/ByteCode/unions.cpp +++ b/clang/test/AST/ByteCode/unions.cpp @@ -79,10 +79,9 @@ namespace DefaultInit { constexpr U1 u1; /// OK. - constexpr int foo() { // expected-error {{never produces a constant expression}} + constexpr int foo() { U1 u; - return u.a; // both-note {{read of member 'a' of union with active member 'b'}} \ - // expected-note {{read of member 'a' of union with active member 'b'}} + return u.a; // both-note {{read of member 'a' of union with active member 'b'}} } static_assert(foo() == 42); // both-error {{not an integral constant expression}} \ // both-note {{in call to}} @@ -861,6 +860,73 @@ namespace CopyCtorMutable { // both-note {{in call}} } + +namespace NonTrivialCtor { + struct A { int x = 1; constexpr int f() { return 1; } }; + struct B : A { int y = 1; constexpr int g() { return 2; } }; + struct C { + int x; + constexpr virtual int f() = 0; + }; + struct D : C { + int y; + constexpr virtual int f() override { return 3; } + }; + + union U { + int n; + B b; + D d; + }; + + consteval int test(int which) { + if (which == 0) {} + + U u{.n = 5}; + assert_active(u); + assert_active(u.n); + assert_inactive(u.b); + + switch (which) { + case 0: + u.b.x = 10; // both-note {{assignment to member 'b' of union with active member 'n'}} + return u.b.f(); + case 1: + u.b.y = 10; // both-note {{assignment to member 'b' of union with active member 'n'}} + return u.b.g(); + case 2: + u.d.x = 10; // both-note {{assignment to member 'd' of union with active member 'n'}} + return u.d.f(); + case 3: + u.d.y = 10; // both-note {{assignment to member 'd' of union with active member 'n'}} + return u.d.f(); + } + + return 1; + } + static_assert(test(0)); // both-error {{not an integral constant expression}} \ + // both-note {{in call}} + static_assert(test(1)); // both-error {{not an integral constant expression}} \ + // both-note {{in call}} + static_assert(test(2)); // both-error {{not an integral constant expression}} \ + // both-note {{in call}} + static_assert(test(3)); // both-error {{not an integral constant expression}} \ + // both-note {{in call}} + +} + +namespace PrimitiveFieldInitActivates { + /// The initializer of a needs the field to be active _before_ it's visited. + template<int> struct X {}; + union V { + int a, b; + constexpr V(X<0>) : a(a = 1) {} // ok + constexpr V(X<2>) : a() { b = 1; } // ok + }; + constinit V v0 = X<0>(); + constinit V v2 = X<2>(); +} + #endif namespace AddressComparison { diff --git a/clang/test/AST/ast-dump-APValue-lvalue.cpp b/clang/test/AST/ast-dump-APValue-lvalue.cpp index 51d22a5..f4cf2f5 100644 --- a/clang/test/AST/ast-dump-APValue-lvalue.cpp +++ b/clang/test/AST/ast-dump-APValue-lvalue.cpp @@ -67,6 +67,10 @@ void Test(int (&arr)[10]) { // CHECK-NEXT: | |-value: LValue Base=TypeInfoLValue typeid(int), Null=0, Offset=0, HasPath=1, PathLength=0, Path=() constexpr int(MP::*pmi) = (int MP::*)&P::x; - // CHECK: `-VarDecl {{.*}} <col:{{.*}}, col:{{.*}}> col:{{.*}} pmi 'int (MP::*const)' constexpr cinit - // CHECK-NEXT: |-value: MemberPointer MP::x + // CHECK: | `-VarDecl {{.*}} <col:{{.*}}, col:{{.*}}> col:{{.*}} pmi 'int (MP::*const)' constexpr cinit + // CHECK-NEXT: | |-value: MemberPointer MP::x + + constexpr int(MP::*pmn) = (int MP::*)nullptr; + // CHECK: `-VarDecl {{.*}} <col:{{.*}}, col:{{.*}}> col:{{.*}} pmn 'int (MP::*const)' constexpr cinit + // CHECK-NEXT: |-value: MemberPointer null } diff --git a/clang/test/Analysis/malloc.c b/clang/test/Analysis/malloc.c index 27a04ff..a9828cf 100644 --- a/clang/test/Analysis/malloc.c +++ b/clang/test/Analysis/malloc.c @@ -1954,9 +1954,23 @@ int conjure(void); void testExtent(void) { int x = conjure(); clang_analyzer_dump(x); - // expected-warning-re@-1 {{{{^conj_\$[[:digit:]]+{int, LC1, S[[:digit:]]+, #1}}}}}} + // expected-warning-re@-1 {{{{^conj_\$[[:digit:]]+{int, LC[[:digit:]]+, S[[:digit:]]+, #1}}}}}} int *p = (int *)malloc(x); clang_analyzer_dumpExtent(p); - // expected-warning-re@-1 {{{{^conj_\$[[:digit:]]+{int, LC1, S[[:digit:]]+, #1}}}}}} + // expected-warning-re@-1 {{{{^conj_\$[[:digit:]]+{int, LC[[:digit:]]+, S[[:digit:]]+, #1}}}}}} free(p); } + +void gh149754(void *p) { + // This testcase demonstrates an unusual situation where a certain symbol + // (the value of `p`) is released (more precisely, transitions from + // untracked state to Released state) twice within the same bug path because + // the `EvalAssume` callback resets it to untracked state after the first + // time when it is released. This caused the failure of an assertion, which + // was since then removed for the codebase. + if (!realloc(p, 8)) { + realloc(p, 8); + free(p); // expected-warning {{Attempt to free released memory}} + } + // expected-warning@+1 {{Potential memory leak}} +} diff --git a/clang/test/CIR/CodeGen/array-ctor.cpp b/clang/test/CIR/CodeGen/array-ctor.cpp new file mode 100644 index 0000000..b3d81a8 --- /dev/null +++ b/clang/test/CIR/CodeGen/array-ctor.cpp @@ -0,0 +1,106 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-cir -mmlir --mlir-print-ir-before=cir-lowering-prepare %s -o - 2>&1 | FileCheck --check-prefixes=CIR-BEFORE-LPP %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +struct S { + S(); +}; + +void foo() { + S s[42]; +} + +// CIR-BEFORE-LPP: cir.func dso_local @_Z3foov() +// CIR-BEFORE-LPP: %[[ARRAY:.*]] = cir.alloca !cir.array<!rec_S x 42>, !cir.ptr<!cir.array<!rec_S x 42>>, ["s", init] +// CIR-BEFORE-LPP: cir.array.ctor %[[ARRAY]] : !cir.ptr<!cir.array<!rec_S x 42>> { +// CIR-BEFORE-LPP: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_S>): +// CIR-BEFORE-LPP: cir.call @_ZN1SC1Ev(%[[ARG]]) : (!cir.ptr<!rec_S>) -> () +// CIR-BEFORE-LPP: cir.yield +// CIR-BEFORE-LPP: } +// CIR-BEFORE-LPP: cir.return +// CIR-BEFORE-LPP: } + +// CIR: cir.func dso_local @_Z3foov() +// CIR: %[[ARRAY:.*]] = cir.alloca !cir.array<!rec_S x 42>, !cir.ptr<!cir.array<!rec_S x 42>>, ["s", init] +// CIR: %[[CONST42:.*]] = cir.const #cir.int<42> : !u64i +// CIR: %[[DECAY:.*]] = cir.cast(array_to_ptrdecay, %[[ARRAY]] : !cir.ptr<!cir.array<!rec_S x 42>>), !cir.ptr<!rec_S> +// CIR: %[[END_PTR:.*]] = cir.ptr_stride(%[[DECAY]] : !cir.ptr<!rec_S>, %[[CONST42]] : !u64i), !cir.ptr<!rec_S> +// CIR: %[[ITER:.*]] = cir.alloca !cir.ptr<!rec_S>, !cir.ptr<!cir.ptr<!rec_S>>, ["__array_idx"] +// CIR: cir.store %[[DECAY]], %[[ITER]] : !cir.ptr<!rec_S>, !cir.ptr<!cir.ptr<!rec_S>> +// CIR: cir.do { +// CIR: %[[CURRENT:.*]] = cir.load %[[ITER]] : !cir.ptr<!cir.ptr<!rec_S>>, !cir.ptr<!rec_S> +// CIR: %[[CONST1:.*]] = cir.const #cir.int<1> : !u64i +// CIR: cir.call @_ZN1SC1Ev(%[[CURRENT]]) : (!cir.ptr<!rec_S>) -> () +// CIR: %[[NEXT:.*]] = cir.ptr_stride(%[[CURRENT]] : !cir.ptr<!rec_S>, %[[CONST1]] : !u64i), !cir.ptr<!rec_S> +// CIR: cir.store %[[NEXT]], %[[ITER]] : !cir.ptr<!rec_S>, !cir.ptr<!cir.ptr<!rec_S>> +// CIR: cir.yield +// CIR: } while { +// CIR: %[[CURRENT2:.*]] = cir.load %[[ITER]] : !cir.ptr<!cir.ptr<!rec_S>>, !cir.ptr<!rec_S> +// CIR: %[[CMP:.*]] = cir.cmp(ne, %[[CURRENT2]], %[[END_PTR]]) : !cir.ptr<!rec_S>, !cir.bool +// CIR: cir.condition(%[[CMP]]) +// CIR: } +// CIR: cir.return +// CIR: } + +// LLVM: define dso_local void @_Z3foov() +// LLVM: %[[ARRAY:.*]] = alloca [42 x %struct.S] +// LLVM: %[[START:.*]] = getelementptr %struct.S, ptr %[[ARRAY]], i32 0 +// LLVM: %[[END:.*]] = getelementptr %struct.S, ptr %[[START]], i64 42 +// LLVM: %[[ITER:.*]] = alloca ptr +// LLVM: store ptr %[[START]], ptr %[[ITER]] +// LLVM: br label %[[LOOP:.*]] +// LLVM: [[COND:.*]]: +// LLVM: %[[CURRENT_CHECK:.*]] = load ptr, ptr %[[ITER]] +// LLVM: %[[DONE:.*]] = icmp ne ptr %[[CURRENT_CHECK]], %[[END]] +// LLVM: br i1 %[[DONE]], label %[[LOOP]], label %[[EXIT:.*]] +// LLVM: [[LOOP]]: +// LLVM: %[[CURRENT:.*]] = load ptr, ptr %[[ITER]] +// LLVM: call void @_ZN1SC1Ev(ptr %[[CURRENT]]) +// LLVM: %[[NEXT:.*]] = getelementptr %struct.S, ptr %[[CURRENT]], i64 1 +// LLVM: store ptr %[[NEXT]], ptr %[[ITER]] +// LLVM: br label %[[COND]] +// LLVM: [[EXIT]]: +// LLVM: ret void + +// OGCG: define dso_local void @_Z3foov() +// OGCG: %[[ARRAY:.*]] = alloca [42 x %struct.S] +// OGCG: %[[START:.*]] = getelementptr{{.*}} %struct.S{{.*}} +// OGCG: %[[END:.*]] = getelementptr{{.*}} %struct.S{{.*}} i64 42 +// OGCG: br label %[[LOOP:.*]] +// OGCG: [[LOOP]]: +// OGCG: %[[CURRENT:.*]] = phi ptr [ %[[START]], %{{.*}} ], [ %[[NEXT:.*]], %[[LOOP]] ] +// OGCG: call void @_ZN1SC1Ev(ptr{{.*}}) +// OGCG: %[[NEXT]] = getelementptr{{.*}} %struct.S{{.*}} i64 1 +// OGCG: %[[DONE:.*]] = icmp eq ptr %[[NEXT]], %[[END]] +// OGCG: br i1 %[[DONE]], label %[[EXIT:.*]], label %[[LOOP]] +// OGCG: [[EXIT]]: +// OGCG: ret void + +void zero_sized() { + S s[0]; +} + +// CIR-BEFORE-LPP: cir.func dso_local @_Z10zero_sizedv() +// CIR-BEFORE-LPP: cir.alloca !cir.array<!rec_S x 0>, !cir.ptr<!cir.array<!rec_S x 0>>, ["s"] +// CIR-BEFORE-LPP-NOT: cir.array.ctor +// CIR-BEFORE-LPP: cir.return + +// CIR: cir.func dso_local @_Z10zero_sizedv() +// CIR: cir.alloca !cir.array<!rec_S x 0>, !cir.ptr<!cir.array<!rec_S x 0>>, ["s"] +// CIR-NOT: cir.do +// CIR-NOT: cir.call @_ZN1SC1Ev +// CIR: cir.return + +// LLVM: define dso_local void @_Z10zero_sizedv() +// LLVM: alloca [0 x %struct.S] +// LLVM-NOT: call void @_ZN1SC1Ev +// LLVM: ret void + +// OGCG: define dso_local void @_Z10zero_sizedv() +// OGCG: alloca [0 x %struct.S] +// OGCG-NOT: call void @_ZN1SC1Ev +// OGCG: ret void diff --git a/clang/test/CIR/CodeGen/bitfields.c b/clang/test/CIR/CodeGen/bitfields.c index a73c076..869a7c9 100644 --- a/clang/test/CIR/CodeGen/bitfields.c +++ b/clang/test/CIR/CodeGen/bitfields.c @@ -315,3 +315,51 @@ void unOp(S* s) { // OGCG: [[TMP12:%.*]] = shl i64 [[TMP8]], 62 // OGCG: [[TMP13:%.*]] = ashr i64 [[TMP12]], 62 // OGCG: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32 + +void binOp(S* s) { + s->d |= 42; +} + +// CIR: cir.func {{.*@binOp}} +// CIR: [[TMP0:%.*]] = cir.const #cir.int<42> : !s32i +// CIR: [[TMP1:%.*]] = cir.get_member {{.*}}[0] {name = "d"} : !cir.ptr<!rec_S> -> !cir.ptr<!u64i> +// CIR: [[TMP2:%.*]] = cir.get_bitfield align(4) (#bfi_d, [[TMP1]] : !cir.ptr<!u64i>) -> !s32i +// CIR: [[TMP3:%.*]] = cir.binop(or, [[TMP2]], [[TMP0]]) : !s32i +// CIR: cir.set_bitfield align(4) (#bfi_d, [[TMP1]] : !cir.ptr<!u64i>, [[TMP3]] : !s32i) + +// LLVM: define {{.*@binOp}} +// LLVM: [[TMP0:%.*]] = load ptr, ptr {{.*}}, align 8 +// LLVM: [[TMP1:%.*]] = getelementptr %struct.S, ptr [[TMP0]], i32 0, i32 0 +// LLVM: [[TMP2:%.*]] = load i64, ptr [[TMP1]], align 4 +// LLVM: [[TMP3:%.*]] = shl i64 [[TMP2]], 13 +// LLVM: [[TMP4:%.*]] = ashr i64 [[TMP3]], 62 +// LLVM: [[TMP5:%.*]] = trunc i64 [[TMP4]] to i32 +// LLVM: [[TMP6:%.*]] = or i32 [[TMP5]], 42 +// LLVM: [[TMP7:%.*]] = zext i32 [[TMP6]] to i64 +// LLVM: [[TMP8:%.*]] = load i64, ptr [[TMP1]], align 4 +// LLVM: [[TMP9:%.*]] = and i64 [[TMP7]], 3 +// LLVM: [[TMP10:%.*]] = shl i64 [[TMP9]], 49 +// LLVM: [[TMP11:%.*]] = and i64 [[TMP8]], -1688849860263937 +// LLVM: [[TMP12:%.*]] = or i64 [[TMP11]], [[TMP10]] +// LLVM: store i64 [[TMP12]], ptr [[TMP1]], align 4 +// LLVM: [[TMP13:%.*]] = shl i64 [[TMP9]], 62 +// LLVM: [[TMP14:%.*]] = ashr i64 [[TMP13]], 62 +// LLVM: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32 + +// OGCG: define {{.*@binOp}} +// OGCG: [[TMP0:%.*]] = load ptr, ptr %s.addr, align 8 +// OGCG: [[TMP1:%.*]] = load i64, ptr [[TMP0]], align 4 +// OGCG: [[TMP2:%.*]] = shl i64 [[TMP1]], 13 +// OGCG: [[TMP3:%.*]] = ashr i64 [[TMP2]], 62 +// OGCG: [[TMP4:%.*]] = trunc i64 [[TMP3]] to i32 +// OGCG: [[TMP5:%.*]] = or i32 [[TMP4]], 42 +// OGCG: [[TMP6:%.*]] = zext i32 [[TMP5]] to i64 +// OGCG: [[TMP7:%.*]] = load i64, ptr [[TMP0]], align 4 +// OGCG: [[TMP8:%.*]] = and i64 [[TMP6]], 3 +// OGCG: [[TMP9:%.*]] = shl i64 [[TMP8]], 49 +// OGCG: [[TMP10:%.*]] = and i64 [[TMP7]], -1688849860263937 +// OGCG: [[TMP11:%.*]] = or i64 [[TMP10]], [[TMP9]] +// OGCG: store i64 [[TMP11]], ptr [[TMP0]], align 4 +// OGCG: [[TMP12:%.*]] = shl i64 [[TMP8]], 62 +// OGCG: [[TMP13:%.*]] = ashr i64 [[TMP12]], 62 +// OGCG: [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32 diff --git a/clang/test/CIR/CodeGen/cleanup.cpp b/clang/test/CIR/CodeGen/cleanup.cpp new file mode 100644 index 0000000..41961513 --- /dev/null +++ b/clang/test/CIR/CodeGen/cleanup.cpp @@ -0,0 +1,83 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s + +struct Struk { + ~Struk(); +}; + +// CHECK: !rec_Struk = !cir.record<struct "Struk" padded {!u8i}> + +// CHECK: cir.func{{.*}} @_ZN5StrukD1Ev(!cir.ptr<!rec_Struk>) + +void test_cleanup() { + Struk s; +} + +// CHECK: cir.func{{.*}} @_Z12test_cleanupv() +// CHECK: %[[S_ADDR:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["s"] +// CHECK: cir.call @_ZN5StrukD1Ev(%[[S_ADDR]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: cir.return + +void test_cleanup_ifelse(bool b) { + if (b) { + Struk s; + } else { + Struk s; + } +} + +// CHECK: cir.func{{.*}} @_Z19test_cleanup_ifelseb(%arg0: !cir.bool +// CHECK: cir.scope { +// CHECK: %[[B:.*]] = cir.load{{.*}} %0 : !cir.ptr<!cir.bool> +// CHECK: cir.if %[[B]] { +// CHECK: %[[S:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["s"] +// CHECK: cir.call @_ZN5StrukD1Ev(%[[S]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: } else { +// CHECK: %[[S_TOO:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["s"] +// CHECK: cir.call @_ZN5StrukD1Ev(%[[S_TOO]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: } +// CHECK: } +// CHECK: cir.return + +void test_cleanup_for() { + for (int i = 0; i < 10; i++) { + Struk s; + } +} + +// CHECK: cir.func{{.*}} @_Z16test_cleanup_forv() +// CHECK: cir.scope { +// CHECK: cir.for : cond { +// CHECK: } body { +// CHECK: cir.scope { +// CHECK: %[[S:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["s"] +// CHECK: cir.call @_ZN5StrukD1Ev(%[[S]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: } +// CHECK: cir.yield +// CHECK: } step { +// CHECK: } +// CHECK: } +// CHECK: cir.return + +void test_cleanup_nested() { + Struk outer; + { + Struk middle; + { + Struk inner; + } + } +} + +// CHECK: cir.func{{.*}} @_Z19test_cleanup_nestedv() +// CHECK: %[[OUTER:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["outer"] +// CHECK: cir.scope { +// CHECK: %[[MIDDLE:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["middle"] +// CHECK: cir.scope { +// CHECK: %[[INNER:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["inner"] +// CHECK: cir.call @_ZN5StrukD1Ev(%[[INNER]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: } +// CHECK: cir.call @_ZN5StrukD1Ev(%[[MIDDLE]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: } +// CHECK: cir.call @_ZN5StrukD1Ev(%[[OUTER]]) nothrow : (!cir.ptr<!rec_Struk>) -> () +// CHECK: cir.return diff --git a/clang/test/CIR/CodeGen/complex-cast.cpp b/clang/test/CIR/CodeGen/complex-cast.cpp new file mode 100644 index 0000000..0881057 --- /dev/null +++ b/clang/test/CIR/CodeGen/complex-cast.cpp @@ -0,0 +1,358 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir -mmlir --mlir-print-ir-before=cir-canonicalize -o %t.cir %s 2>&1 | FileCheck --check-prefix=CIR-BEFORE %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir -mmlir --mlir-print-ir-after=cir-lowering-prepare -o %t.cir %s 2>&1 | FileCheck --check-prefixes=CIR-AFTER %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +double _Complex cd; +float _Complex cf; +int _Complex ci; +short _Complex cs; +double sd; +int si; +bool b; + +void scalar_to_complex() { + cd = sd; + ci = si; + cd = si; + ci = sd; +} + +// CIR-BEFORE: %[[FP_TO_COMPLEX:.*]] = cir.cast(float_to_complex, %{{.*}} : !cir.double), !cir.complex<!cir.double> + +// CIR-AFTER: %[[REAL:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.double>, !cir.double +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.fp<0.000000e+00> : !cir.double +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !cir.double -> !cir.complex<!cir.double> + +// LLVM: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// LLVM-NEXT: %[[TMP:.*]] = insertvalue { double, double } undef, double %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { double, double } %[[TMP]], double 0.000000e+00, 1 +// LLVM-NEXT: store { double, double } %[[COMPLEX]], ptr {{.*}}, align 8 + +// OGCG: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: store double %[[REAL]], ptr {{.*}}, align 8 +// OGCG: store double 0.000000e+00, ptr getelementptr inbounds nuw ({ double, double }, ptr @cd, i32 0, i32 1), align 8 + +// CIR-BEFORE: %[[INT_TO_COMPLEX:.*]] = cir.cast(int_to_complex, %{{.*}} : !s32i), !cir.complex<!s32i> + +// CIR-AFTER: %[[REAL:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!s32i>, !s32i +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.int<0> : !s32i +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !s32i -> !cir.complex<!s32i> + +// LLVM: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// LLVM-NEXT: %[[TMP:.*]] = insertvalue { i32, i32 } undef, i32 %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { i32, i32 } %[[TMP]], i32 0, 1 +// LLVM-NEXT: store { i32, i32 } %[[COMPLEX]], ptr {{.*}}, align 4 + +// OGCG: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: store i32 %[[REAL]], ptr {{.*}}, align 4 +// OGCG: store i32 0, ptr getelementptr inbounds nuw ({ i32, i32 }, ptr @ci, i32 0, i32 1), align 4 + +// CIR-BEFORE: %[[INT_TO_FP:.*]] = cir.cast(int_to_float, %{{.*}} : !s32i), !cir.double +// CIR-BEFORE: %[[FP_TO_COMPLEX:.*]] = cir.cast(float_to_complex, %[[INT_TO_FP]] : !cir.double), !cir.complex<!cir.double> + +// CIR-AFTER: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!s32i>, !s32i +// CIR-AFTER-NEXT: %[[REAL:.*]] = cir.cast(int_to_float, %[[TMP]] : !s32i), !cir.double +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.fp<0.000000e+00> : !cir.double +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !cir.double -> !cir.complex<!cir.double> + +// LLVM: %[[TMP:.*]] = load i32, ptr {{.*}}, align 4 +// LLVM-NEXT: %[[REAL:.*]] = sitofp i32 %[[TMP]] to double +// LLVM-NEXT: %[[TMP_2:.*]] = insertvalue { double, double } undef, double %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { double, double } %[[TMP_2]], double 0.000000e+00, 1 +// LLVM-NEXT: store { double, double } %[[COMPLEX]], ptr {{.*}}, align 8 + +// OGCG: %[[TMP:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: %[[REAL:.*]] = sitofp i32 %[[TMP]] to double +// OGCG: store double %[[REAL]], ptr {{.*}}, align 8 +// OGCG: store double 0.000000e+00, ptr getelementptr inbounds nuw ({ double, double }, ptr {{.*}}, i32 0, i32 1), align 8 + +// CIR-BEFORE: %[[FP_TO_INT:.*]] = cir.cast(float_to_int, %{{.*}} : !cir.double), !s32i +// CIR-BEFORE: %[[INT_TO_COMPLEX:.*]] = cir.cast(int_to_complex, %[[FP_TO_INT]] : !s32i), !cir.complex<!s32i> + +// CIR-AFTER: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.double>, !cir.double +// CIR-AFTER-NEXT: %[[REAL:.*]] = cir.cast(float_to_int, %[[TMP]] : !cir.double), !s32i +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.int<0> : !s32i +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !s32i -> !cir.complex<!s32i> + +// LLVM: %[[TMP:.*]] = load double, ptr {{.*}}, align 8 +// LLVM-NEXT: %[[REAL:.*]] = fptosi double %[[TMP]] to i32 +// LLVM-NEXT: %[[TMP_2:.*]] = insertvalue { i32, i32 } undef, i32 %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { i32, i32 } %[[TMP_2]], i32 0, 1 +// LLVM-NEXT: store { i32, i32 } %[[COMPLEX]], ptr {{.*}}, align 4 + +// OGCG: %[[TMP:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: %[[REAL:.*]] = fptosi double %[[TMP]] to i32 +// OGCG: store i32 %[[REAL]], ptr {{.*}}, align 4 +// OGCG: store i32 0, ptr getelementptr inbounds nuw ({ i32, i32 }, ptr {{.*}}, i32 0, i32 1), align 4 + +void scalar_to_complex_explicit() { + cd = (double _Complex)sd; + ci = (int _Complex)si; + cd = (double _Complex)si; + ci = (int _Complex)sd; +} + +// CIR-BEFORE: %[[FP_TO_COMPLEX:.*]] = cir.cast(float_to_complex, %{{.*}} : !cir.double), !cir.complex<!cir.double> + +// CIR-AFTER: %[[REAL:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.double>, !cir.double +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.fp<0.000000e+00> : !cir.double +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !cir.double -> !cir.complex<!cir.double> + +// LLVM: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// LLVM-NEXT: %[[TMP:.*]] = insertvalue { double, double } undef, double %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { double, double } %[[TMP]], double 0.000000e+00, 1 +// LLVM-NEXT: store { double, double } %[[COMPLEX]], ptr {{.*}}, align 8 + +// OGCG: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: store double %[[REAL]], ptr {{.*}}, align 8 +// OGCG: store double 0.000000e+00, ptr getelementptr inbounds nuw ({ double, double }, ptr @cd, i32 0, i32 1), align 8 + +// CIR-BEFORE: %[[INT_TO_COMPLEX:.*]] = cir.cast(int_to_complex, %{{.*}} : !s32i), !cir.complex<!s32i> + +// CIR-AFTER: %[[REAL:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!s32i>, !s32i +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.int<0> : !s32i +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !s32i -> !cir.complex<!s32i> + +// LLVM: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// LLVM-NEXT: %[[TMP:.*]] = insertvalue { i32, i32 } undef, i32 %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { i32, i32 } %[[TMP]], i32 0, 1 +// LLVM-NEXT: store { i32, i32 } %[[COMPLEX]], ptr {{.*}}, align 4 + +// OGCG: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: store i32 %[[REAL]], ptr {{.*}}, align 4 +// OGCG: store i32 0, ptr getelementptr inbounds nuw ({ i32, i32 }, ptr @ci, i32 0, i32 1), align 4 + +// CIR-BEFORE: %[[INT_TO_FP:.*]] = cir.cast(int_to_float, %{{.*}} : !s32i), !cir.double +// CIR-BEFORE: %[[FP_TO_COMPLEX:.*]] = cir.cast(float_to_complex, %[[INT_TO_FP]] : !cir.double), !cir.complex<!cir.double> + +// CIR-AFTER: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!s32i>, !s32i +// CIR-AFTER-NEXT: %[[REAL:.*]] = cir.cast(int_to_float, %[[TMP]] : !s32i), !cir.double +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.fp<0.000000e+00> : !cir.double +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !cir.double -> !cir.complex<!cir.double> + +// LLVM: %[[TMP:.*]] = load i32, ptr {{.*}}, align 4 +// LLVM-NEXT: %[[REAL:.*]] = sitofp i32 %[[TMP]] to double +// LLVM-NEXT: %[[TMP_2:.*]] = insertvalue { double, double } undef, double %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { double, double } %[[TMP_2]], double 0.000000e+00, 1 +// LLVM-NEXT: store { double, double } %[[COMPLEX]], ptr {{.*}}, align 8 + +// OGCG: %[[TMP:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: %[[REAL:.*]] = sitofp i32 %[[TMP]] to double +// OGCG: store double %[[REAL]], ptr {{.*}}, align 8 +// OGCG: store double 0.000000e+00, ptr getelementptr inbounds nuw ({ double, double }, ptr {{.*}}, i32 0, i32 1), align 8 + +// CIR-BEFORE: %[[FP_TO_INT:.*]] = cir.cast(float_to_int, %{{.*}} : !cir.double), !s32i +// CIR-BEFORE: %[[INT_TO_COMPLEX:.*]] = cir.cast(int_to_complex, %[[FP_TO_INT]] : !s32i), !cir.complex<!s32i> + +// CIR-AFTER: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.double>, !cir.double +// CIR-AFTER-NEXT: %[[REAL:.*]] = cir.cast(float_to_int, %[[TMP]] : !cir.double), !s32i +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.const #cir.int<0> : !s32i +// CIR-AFTER-NEXT: %{{.*}} = cir.complex.create %[[REAL]], %[[IMAG]] : !s32i -> !cir.complex<!s32i> + +// LLVM: %[[TMP:.*]] = load double, ptr {{.*}}, align 8 +// LLVM-NEXT: %[[REAL:.*]] = fptosi double %[[TMP]] to i32 +// LLVM-NEXT: %[[TMP_2:.*]] = insertvalue { i32, i32 } undef, i32 %[[REAL]], 0 +// LLVM-NEXT: %[[COMPLEX:.*]] = insertvalue { i32, i32 } %[[TMP_2]], i32 0, 1 +// LLVM-NEXT: store { i32, i32 } %[[COMPLEX]], ptr {{.*}}, align 4 + +// OGCG: %[[TMP:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: %[[REAL:.*]] = fptosi double %[[TMP]] to i32 +// OGCG: store i32 %[[REAL]], ptr {{.*}}, align 4 +// OGCG: store i32 0, ptr getelementptr inbounds nuw ({ i32, i32 }, ptr {{.*}}, i32 0, i32 1), align 4 + +void complex_to_scalar() { + sd = (double)cd; + si = (int)ci; + sd = (double)ci; + si = (int)cd; +} + +// CIR-BEFORE: %[[FP_TO_COMPLEX_REAL:.*]] = cir.cast(float_complex_to_real, %{{.*}} : !cir.complex<!cir.double>), !cir.double + +// CIR-AFTER: %{{.*}} = cir.complex.real %{{.*}} : !cir.complex<!cir.double> -> !cir.double + +// LLVM: %[[REAL:.*]] = extractvalue { double, double } %{{.*}}, 0 +// LLVM: store double %[[REAL]], ptr {{.*}}, align 8 + +// OGCG: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: store double %[[REAL]], ptr {{.*}}, align 8 + +// CIR-BEFORE: %[[INT_COMPLEX_TO_REAL:.*]] = cir.cast(int_complex_to_real, %{{.*}} : !cir.complex<!s32i>), !s32i + +// CIR-AFTER: %{{.*}} = cir.complex.real %{{.*}} : !cir.complex<!s32i> -> !s32i + +// LLVM: %[[REAL:.*]] = extractvalue { i32, i32 } %{{.*}}, 0 +// LLVM: store i32 %[[REAL]], ptr {{.*}}, align 4 + +// OGCG: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: store i32 %[[REAL]], ptr {{.*}}, align 4 + +// CIR-BEFORE: %[[INT_COMPLEX_TO_REAL:.*]] = cir.cast(int_complex_to_real, %{{.*}} : !cir.complex<!s32i>), !s32i +// CIR-BEFORE: %[[INT_TO_FP:.*]] = cir.cast(int_to_float, %[[INT_COMPLEX_TO_REAL]] : !s32i), !cir.double + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!s32i> -> !s32i +// CIR-AFTER-NEXT: %{{.*}} = cir.cast(int_to_float, %[[REAL]] : !s32i), !cir.double + +// LLVM: %[[REAL:.*]] = extractvalue { i32, i32 } %{{.+}}, 0 +// LLVM-NEXT: %[[REAL_TO_DOUBLE:.*]] = sitofp i32 %[[REAL]] to double +// LLVM-NEXT: store double %[[REAL_TO_DOUBLE]], ptr {{.*}}, align 8 + +// OGCG: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: %[[INT_TO_FP:.*]] = sitofp i32 %[[REAL]] to double +// OGCG: store double %[[INT_TO_FP]], ptr {{.*}}, align 8 + +// CIR-BEFORE: %[[FP_TO_COMPLEX_REAL:.*]] = cir.cast(float_complex_to_real, %{{.*}} : !cir.complex<!cir.double>), !cir.double +// CIR-BEFORE: %[[FP_TO_INT:.*]] = cir.cast(float_to_int, %[[FP_TO_COMPLEX_REAL]] : !cir.double), !s32i + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!cir.double> -> !cir.double +// CIR-AFTER-NEXT: %{{.*}} = cir.cast(float_to_int, %[[REAL]] : !cir.double), !s32i + +// LLVM: %[[REAL:.*]] = extractvalue { double, double } %{{.+}}, 0 +// LLVM-NEXT: %[[REAL_TO_INT:.*]] = fptosi double %[[REAL]] to i32 +// LLVM-NEXT: store i32 %[[REAL_TO_INT]], ptr {{.*}}, align 4 + +// OGCG: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: %[[FP_TO_INT:.*]] = fptosi double %[[REAL]] to i32 +// OGCG: store i32 %[[FP_TO_INT]], ptr {{.*}}, align 4 + +void complex_to_bool() { + b = (bool)cd; + b = (bool)ci; +} + +// CIR-BEFORE: %[[FP_COMPLEX_TO_BOOL:.*]] = cir.cast(float_complex_to_bool, %{{.*}} : !cir.complex<!cir.double>), !cir.bool + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!cir.double> -> !cir.double +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.complex.imag %{{.*}} : !cir.complex<!cir.double> -> !cir.double +// CIR-AFTER-NEXT: %[[REAL_TO_BOOL:.*]] = cir.cast(float_to_bool, %[[REAL]] : !cir.double), !cir.bool +// CIR-AFTER-NEXT: %[[IMAG_TO_BOOL:.*]] = cir.cast(float_to_bool, %[[IMAG]] : !cir.double), !cir.bool +// CIR-AFTER-NEXT: %[[CONST_TRUE:.*]] = cir.const #true +// CIR-AFTER-NEXT: %{{.*}} = cir.select if %[[REAL_TO_BOOL]] then %[[CONST_TRUE]] else %[[IMAG_TO_BOOL]] : (!cir.bool, !cir.bool, !cir.bool) -> !cir.bool + +// LLVM: %[[REAL:.*]] = extractvalue { double, double } %{{.*}}, 0 +// LLVM-NEXT: %[[IMAG:.*]] = extractvalue { double, double } %{{.*}}, 1 +// LLVM-NEXT: %[[REAL_TO_BOOL:.*]] = fcmp une double %[[REAL]], 0.000000e+00 +// LLVM-NEXT: %[[IMAG_TO_BOOL:.*]] = fcmp une double %[[IMAG]], 0.000000e+00 +// LLVM-NEXT: %[[OR:.*]] = or i1 %[[REAL_TO_BOOL]], %[[IMAG_TO_BOOL]] +// LLVM-NEXT: %[[RESULT:.*]] = zext i1 %[[OR]] to i8 +// LLVM-NEXT: store i8 %[[RESULT]], ptr {{.*}}, align 1 + +// OGCG: %[[REAL:.*]] = load double, ptr {{.*}}, align 8 +// OGCG: %[[IMAG:.*]] = load double, ptr getelementptr inbounds nuw ({ double, double }, ptr {{.*}}, i32 0, i32 1), align 8 +// OGCG: %[[REAL_TO_BOOL:.*]] = fcmp une double %[[REAL]], 0.000000e+00 +// OGCG: %[[IMAG_TO_BOOL:.*]] = fcmp une double %[[IMAG]], 0.000000e+00 +// OGCG: %[[COMPLEX_TO_BOOL:.*]] = or i1 %[[REAL_TO_BOOL]], %[[IMAG_TO_BOOL]] +// OGCG: %[[BOOL_TO_INT:.*]] = zext i1 %[[COMPLEX_TO_BOOL]] to i8 +// OGCG: store i8 %[[BOOL_TO_INT]], ptr {{.*}}, align 1 + +// CIR-BEFORE: %[[INT_COMPLEX_TO_BOOL:.*]] = cir.cast(int_complex_to_bool, %{{.*}} : !cir.complex<!s32i>), !cir.bool + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!s32i> -> !s32i +// CIR-AFTER-NEXT: %[[IMAG:.*]] = cir.complex.imag %{{.*}} : !cir.complex<!s32i> -> !s32i +// CIR-AFTER-NEXT: %[[REAL_TO_BOOL:.*]] = cir.cast(int_to_bool, %[[REAL]] : !s32i), !cir.bool +// CIR-AFTER-NEXT: %[[IMAG_TO_BOOL:.*]] = cir.cast(int_to_bool, %[[IMAG]] : !s32i), !cir.bool +// CIR-AFTER-NEXT: %[[CONST_TRUE:.*]] = cir.const #true +// CIR-AFTER-NEXT: %{{.+}} = cir.select if %[[REAL_TO_BOOL]] then %[[CONST_TRUE]] else %[[IMAG_TO_BOOL]] : (!cir.bool, !cir.bool, !cir.bool) -> !cir.bool + +// LLVM: %[[REAL:.*]] = extractvalue { i32, i32 } %{{.*}}, 0 +// LLVM-NEXT: %[[IMAG:.*]] = extractvalue { i32, i32 } %{{.*}}, 1 +// LLVM-NEXT: %[[REAL_TO_BOOL:.*]] = icmp ne i32 %[[REAL]], 0 +// LLVM-NEXT: %[[IMAG_TO_BOOL:.*]] = icmp ne i32 %[[IMAG]], 0 +// LLVM-NEXT: %[[OR:.*]] = or i1 %[[REAL_TO_BOOL]], %[[IMAG_TO_BOOL]] +// LLVM-NEXT: %[[RESULT:.*]] = zext i1 %[[OR]] to i8 +// LLVM-NEXT: store i8 %[[RESULT]], ptr {{.*}}, align 1 + +// OGCG: %[[REAL:.*]] = load i32, ptr {{.*}}, align 4 +// OGCG: %[[IMAG:.*]] = load i32, ptr getelementptr inbounds nuw ({ i32, i32 }, ptr {{.*}}, i32 0, i32 1), align 4 +// OGCG: %[[REAL_TO_BOOL:.*]] = icmp ne i32 %[[REAL]], 0 +// OGCG: %[[IMAG_TO_BOOL:.*]] = icmp ne i32 %[[IMAG]], 0 +// OGCG: %[[COMPLEX_TO_BOOL:.*]] = or i1 %[[REAL_TO_BOOL]], %[[IMAG_TO_BOOL]] +// OGCG: %[[BOOL_TO_INT:.*]] = zext i1 %[[COMPLEX_TO_BOOL]] to i8 +// OGCG: store i8 %[[BOOL_TO_INT]], ptr {{.*}}, align 1 + +void complex_to_complex_cast() { + cd = cf; + ci = cs; +} + +// CIR-BEFORE: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.complex<!cir.float>>, !cir.complex<!cir.float> +// CIR-BEFORE: %[[FP_COMPLEX:.*]] = cir.cast(float_complex, %[[TMP]] : !cir.complex<!cir.float>), !cir.complex<!cir.double> + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!cir.float> -> !cir.float +// CIR-AFTER: %[[IMAG:.*]] = cir.complex.imag %{{.*}} : !cir.complex<!cir.float> -> !cir.float +// CIR-AFTER: %[[REAL_FP_CAST:.*]] = cir.cast(floating, %[[REAL]] : !cir.float), !cir.double +// CIR-AFTER: %[[IMAG_FP_CAST:.*]] = cir.cast(floating, %[[IMAG]] : !cir.float), !cir.double +// CIR-AFTER: %{{.*}} = cir.complex.create %[[REAL_FP_CAST]], %[[IMAG_FP_CAST]] : !cir.double -> !cir.complex<!cir.double> + +// LLVM: %[[REAL:.*]] = extractvalue { float, float } %{{.*}}, 0 +// LLVM: %[[IMAG:.*]] = extractvalue { float, float } %{{.*}}, 1 +// LLVM: %[[REAL_FP_CAST:.*]] = fpext float %[[REAL]] to double +// LLVM: %[[IMAG_FP_CAST:.*]] = fpext float %[[IMAG]] to double +// LLVM: %[[TMP:.*]] = insertvalue { double, double } undef, double %[[REAL_FP_CAST]], 0 +// LLVM: %[[COMPLEX:.*]] = insertvalue { double, double } %[[TMP]], double %[[IMAG_FP_CAST]], 1 +// LLVM: store { double, double } %[[COMPLEX]], ptr {{.*}}, align 8 + +// OGCG: %[[REAL:.*]] = load float, ptr {{.*}}, align 4 +// OGCG: %[[IMAG:.*]] = load float, ptr getelementptr inbounds nuw ({ float, float }, ptr {{.*}}, i32 0, i32 1), align 4 +// OGCG: %[[REAL_FP_CAST:.*]] = fpext float %[[REAL]] to double +// OGCG: %[[IMAG_FP_CAST:.*]] = fpext float %[[IMAG]] to double +// OGCG: store double %[[REAL_FP_CAST]], ptr {{.*}}, align 8 +// OGCG: store double %[[IMAG_FP_CAST]], ptr getelementptr inbounds nuw ({ double, double }, ptr {{.*}}, i32 0, i32 1), align 8 + +// CIR-BEFORE: %[[TMP:.*]] = cir.load{{.*}} %{{.*}} : !cir.ptr<!cir.complex<!s16i>>, !cir.complex<!s16i> +// CIR-BEFORE: %[[INT_COMPLEX:.*]] = cir.cast(int_complex, %[[TMP]] : !cir.complex<!s16i>), !cir.complex<!s32i> + +// CIR-AFTER: %[[REAL:.*]] = cir.complex.real %{{.*}} : !cir.complex<!s16i> -> !s16i +// CIR-AFTER: %[[IMAG:.*]] = cir.complex.imag %{{.*}} : !cir.complex<!s16i> -> !s16i +// CIR-AFTER: %[[REAL_INT_CAST:.*]] = cir.cast(integral, %[[REAL]] : !s16i), !s32i +// CIR-AFTER: %[[IMAG_INT_CAST:.*]] = cir.cast(integral, %[[IMAG]] : !s16i), !s32i +// CIR-AFTER: %{{.*}} = cir.complex.create %[[REAL_INT_CAST]], %[[IMAG_INT_CAST]] : !s32i -> !cir.complex<!s32i> + +// LLVM: %[[REAL:.*]] = extractvalue { i16, i16 } %{{.*}}, 0 +// LLVM: %[[IMAG:.*]] = extractvalue { i16, i16 } %{{.*}}, 1 +// LLVM: %[[REAL_INT_CAST:.*]] = sext i16 %[[REAL]] to i32 +// LLVM: %[[IMAG_INT_CAST:.*]] = sext i16 %[[IMAG]] to i32 +// LLVM: %[[TMP:.*]] = insertvalue { i32, i32 } undef, i32 %[[REAL_INT_CAST]], 0 +// LLVM: %[[COMPLEX:.*]] = insertvalue { i32, i32 } %[[TMP]], i32 %[[IMAG_INT_CAST]], 1 +// LLVM: store { i32, i32 } %[[COMPLEX]], ptr {{.*}}, align 4 + +// OGCG: %[[REAL:.*]] = load i16, ptr {{.*}}, align 2 +// OGCG: %[[IMAG:.*]] = load i16, ptr getelementptr inbounds nuw ({ i16, i16 }, ptr {{.*}}, i32 0, i32 1), align 2 +// OGCG: %[[REAL_INT_CAST:.*]] = sext i16 %[[REAL]] to i32 +// OGCG: %[[IMAG_INT_CAST:.*]] = sext i16 %[[IMAG]] to i32 +// OGCG: store i32 %[[REAL_INT_CAST]], ptr {{.*}}, align 4 +// OGCG: store i32 %[[IMAG_INT_CAST]], ptr getelementptr inbounds nuw ({ i32, i32 }, ptr {{.*}}, i32 0, i32 1), align 4 + +struct CX { + double real; + double imag; +}; + +void lvalue_to_rvalue_bitcast() { + CX a; + double _Complex b = __builtin_bit_cast(double _Complex, a); +} + + +// CIR-BEFORE: %{{.*}} = cir.cast(bitcast, %{{.*}} : !cir.ptr<!rec_CX>), !cir.ptr<!cir.complex<!cir.double>> + +// CIR-AFTER: %{{.*}} = cir.cast(bitcast, %{{.*}} : !cir.ptr<!rec_CX>), !cir.ptr<!cir.complex<!cir.double>> + +// LLVM: %[[PTR_ADDR:.*]] = alloca %struct.CX, i64 1, align 8 +// LLVM: %[[COMPLEX_ADDR:.*]] = alloca { double, double }, i64 1, align 8 +// LLVM: %[[PTR_TO_COMPLEX:.*]] = load { double, double }, ptr %[[PTR_ADDR]], align 8 +// LLVM: store { double, double } %[[PTR_TO_COMPLEX]], ptr %[[COMPLEX_ADDR]], align 8 + +// OGCG: %[[A_ADDR:.*]] = alloca %struct.CX, align 8 +// OGCG: %[[B_ADDR:.*]] = alloca { double, double }, align 8 +// OGCG: %[[A_REAL_PTR:.*]] = getelementptr inbounds nuw { double, double }, ptr %[[A_ADDR]], i32 0, i32 0 +// OGCG: %[[A_REAL:.*]] = load double, ptr %[[A_REAL_PTR]], align 8 +// OGCG: %[[A_IMAG_PTR:.*]] = getelementptr inbounds nuw { double, double }, ptr %[[A_ADDR]], i32 0, i32 1 +// OGCG: %[[A_IMAG:.*]] = load double, ptr %[[A_IMAG_PTR]], align 8 +// OGCG: %[[B_REAL_PTR:.*]] = getelementptr inbounds nuw { double, double }, ptr %[[B_ADDR]], i32 0, i32 0 +// OGCG: %[[B_IMAG_PTR:.*]] = getelementptr inbounds nuw { double, double }, ptr %[[B_ADDR]], i32 0, i32 1 +// OGCG: store double %[[A_REAL]], ptr %[[B_REAL_PTR]], align 8 +// OGCG: store double %[[A_IMAG]], ptr %[[B_IMAG_PTR]], align 8 diff --git a/clang/test/CIR/CodeGen/dtor-alias.cpp b/clang/test/CIR/CodeGen/dtor-alias.cpp new file mode 100644 index 0000000..e37ddab --- /dev/null +++ b/clang/test/CIR/CodeGen/dtor-alias.cpp @@ -0,0 +1,75 @@ +// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -mconstructor-aliases -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s +// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -mconstructor-aliases -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -mconstructor-aliases -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s + +struct B { + ~B(); +}; +B::~B() { +} + +// OGCG: @_ZN1BD1Ev = unnamed_addr alias void (ptr), ptr @_ZN1BD2Ev + +// CHECK: cir.func{{.*}} @_ZN1BD2Ev(%arg0: !cir.ptr<!rec_B> +// CHECK: %[[THIS_ADDR:.*]] = cir.alloca !cir.ptr<!rec_B>, !cir.ptr<!cir.ptr<!rec_B>>, ["this", init] +// CHECK: cir.store %arg0, %[[THIS_ADDR]] +// CHECK: %[[THIS:.*]] = cir.load %[[THIS_ADDR]] : !cir.ptr<!cir.ptr<!rec_B>>, !cir.ptr<!rec_B> + +// CHECK: cir.func{{.*}} private dso_local @_ZN1BD1Ev(!cir.ptr<!rec_B>) alias(@_ZN1BD2Ev) + +// LLVM: define{{.*}} @_ZN1BD2Ev(ptr %[[THIS_ARG:.*]]) +// LLVM: %[[THIS_ADDR:.*]] = alloca ptr +// LLVM: store ptr %[[THIS_ARG]], ptr %[[THIS_ADDR]] +// LLVM: %[[THIS:.*]] = load ptr, ptr %[[THIS_ADDR]] + +// This should be an alias, like the similar OGCG alias above, but that's not +// implemented yet. +// LLVM: declare dso_local void @_ZN1BD1Ev(ptr) + +// OGCG: define{{.*}} @_ZN1BD2Ev(ptr{{.*}} %[[THIS_ARG:.*]]) +// OGCG: %[[THIS_ADDR:.*]] = alloca ptr +// OGCG: store ptr %[[THIS_ARG]], ptr %[[THIS_ADDR]] +// OGCG: %[[THIS:.*]] = load ptr, ptr %[[THIS_ADDR]] + +// The destructor in this case is handled by RAUW rather than aliasing. +struct Struk { + ~Struk() {} +}; + +void baz() { + Struk s; +} + +// CHECK: cir.func{{.*}} @_ZN5StrukD2Ev(%arg0: !cir.ptr<!rec_Struk> +// CHECK: %[[THIS_ADDR:.*]] = cir.alloca !cir.ptr<!rec_Struk>, !cir.ptr<!cir.ptr<!rec_Struk>>, ["this", init] +// CHECK: cir.store %arg0, %[[THIS_ADDR]] : !cir.ptr<!rec_Struk>, !cir.ptr<!cir.ptr<!rec_Struk>> +// CHECK: %[[THIS:.*]] = cir.load %[[THIS_ADDR]] : !cir.ptr<!cir.ptr<!rec_Struk>>, !cir.ptr<!rec_Struk> +// CHECK: cir.return + +// CHECK-NOT: cir.func{{.*}} @_ZN5StrukD1Ev + +// CHECK: cir.func{{.*}} @_Z3bazv() +// CHECK: %[[S_ADDR:.*]] = cir.alloca !rec_Struk, !cir.ptr<!rec_Struk>, ["s"] +// CHECK: cir.call @_ZN5StrukD2Ev(%[[S_ADDR]]) nothrow : (!cir.ptr<!rec_Struk>) -> () + +// LLVM: define linkonce_odr void @_ZN5StrukD2Ev(ptr %[[THIS_ARG]]) +// LLVM: %[[THIS_ADDR:.*]] = alloca ptr +// LLVM: store ptr %[[THIS_ARG]], ptr %[[THIS_ADDR]] +// LLVM: %[[THIS:.*]] = load ptr, ptr %[[THIS_ADDR]] + +// LLVM: define{{.*}} void @_Z3bazv() +// LLVM: %[[S_ADDR:.*]] = alloca %struct.Struk +// LLVM: call void @_ZN5StrukD2Ev(ptr{{.*}} %[[S_ADDR]]) + +// This function gets emitted before the destructor in OGCG. +// OGCG: define{{.*}} void @_Z3bazv() +// OGCG: %[[S_ADDR:.*]] = alloca %struct.Struk +// OGCG: call void @_ZN5StrukD2Ev(ptr{{.*}} %[[S_ADDR]]) + +// OGCG: define linkonce_odr void @_ZN5StrukD2Ev(ptr{{.*}} %[[THIS_ARG]]) +// OGCG: %[[THIS_ADDR:.*]] = alloca ptr +// OGCG: store ptr %[[THIS_ARG]], ptr %[[THIS_ADDR]] +// OGCG: %[[THIS:.*]] = load ptr, ptr %[[THIS_ADDR]] diff --git a/clang/test/CIR/IR/array-ctor.cir b/clang/test/CIR/IR/array-ctor.cir new file mode 100644 index 0000000..2378992 --- /dev/null +++ b/clang/test/CIR/IR/array-ctor.cir @@ -0,0 +1,29 @@ + +// RUN: cir-opt %s | FileCheck %s + +!u8i = !cir.int<u, 8> +!rec_S = !cir.record<struct "S" padded {!u8i}> + +module { + cir.func private @_ZN1SC1Ev(!cir.ptr<!rec_S>) + cir.func dso_local @_Z3foov() { + %0 = cir.alloca !cir.array<!rec_S x 42>, !cir.ptr<!cir.array<!rec_S x 42>>, ["s", init] {alignment = 16 : i64} + cir.array.ctor %0 : !cir.ptr<!cir.array<!rec_S x 42>> { + ^bb0(%arg0: !cir.ptr<!rec_S>): + cir.call @_ZN1SC1Ev(%arg0) : (!cir.ptr<!rec_S>) -> () + cir.yield + } + cir.return + } + + // CHECK: cir.func private @_ZN1SC1Ev(!cir.ptr<!rec_S>) + // CHECK: cir.func dso_local @_Z3foov() { + // CHECK: %0 = cir.alloca !cir.array<!rec_S x 42>, !cir.ptr<!cir.array<!rec_S x 42>>, ["s", init] {alignment = 16 : i64} + // CHECK: cir.array.ctor %0 : !cir.ptr<!cir.array<!rec_S x 42>> { + // CHECK: ^bb0(%arg0: !cir.ptr<!rec_S>): + // CHECK: cir.call @_ZN1SC1Ev(%arg0) : (!cir.ptr<!rec_S>) -> () + // CHECK: cir.yield + // CHECK: } + // CHECK: cir.return + // CHECK: } +} diff --git a/clang/test/CodeGen/AArch64/neon-intrinsics.c b/clang/test/CodeGen/AArch64/neon-intrinsics.c index 6304245..035e1ca 100644 --- a/clang/test/CodeGen/AArch64/neon-intrinsics.c +++ b/clang/test/CodeGen/AArch64/neon-intrinsics.c @@ -8585,7 +8585,7 @@ uint32x2_t test_vqshrun_n_s64(int64x2_t a) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <8 x i8> [[A]], <8 x i8> [[VQSHRUN_N3]], <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15> // CHECK-NEXT: ret <16 x i8> [[SHUFFLE_I]] // -int8x16_t test_vqshrun_high_n_s16(int8x8_t a, int16x8_t b) { +uint8x16_t test_vqshrun_high_n_s16(uint8x8_t a, int16x8_t b) { return vqshrun_high_n_s16(a, b, 3); } @@ -8598,7 +8598,7 @@ int8x16_t test_vqshrun_high_n_s16(int8x8_t a, int16x8_t b) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <4 x i16> [[A]], <4 x i16> [[VQSHRUN_N3]], <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> // CHECK-NEXT: ret <8 x i16> [[SHUFFLE_I]] // -int16x8_t test_vqshrun_high_n_s32(int16x4_t a, int32x4_t b) { +uint16x8_t test_vqshrun_high_n_s32(uint16x4_t a, int32x4_t b) { return vqshrun_high_n_s32(a, b, 9); } @@ -8611,7 +8611,7 @@ int16x8_t test_vqshrun_high_n_s32(int16x4_t a, int32x4_t b) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <2 x i32> [[A]], <2 x i32> [[VQSHRUN_N3]], <4 x i32> <i32 0, i32 1, i32 2, i32 3> // CHECK-NEXT: ret <4 x i32> [[SHUFFLE_I]] // -int32x4_t test_vqshrun_high_n_s64(int32x2_t a, int64x2_t b) { +uint32x4_t test_vqshrun_high_n_s64(uint32x2_t a, int64x2_t b) { return vqshrun_high_n_s64(a, b, 19); } @@ -8810,7 +8810,7 @@ uint32x2_t test_vqrshrun_n_s64(int64x2_t a) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <8 x i8> [[A]], <8 x i8> [[VQRSHRUN_N3]], <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15> // CHECK-NEXT: ret <16 x i8> [[SHUFFLE_I]] // -int8x16_t test_vqrshrun_high_n_s16(int8x8_t a, int16x8_t b) { +uint8x16_t test_vqrshrun_high_n_s16(uint8x8_t a, int16x8_t b) { return vqrshrun_high_n_s16(a, b, 3); } @@ -8823,7 +8823,7 @@ int8x16_t test_vqrshrun_high_n_s16(int8x8_t a, int16x8_t b) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <4 x i16> [[A]], <4 x i16> [[VQRSHRUN_N3]], <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> // CHECK-NEXT: ret <8 x i16> [[SHUFFLE_I]] // -int16x8_t test_vqrshrun_high_n_s32(int16x4_t a, int32x4_t b) { +uint16x8_t test_vqrshrun_high_n_s32(uint16x4_t a, int32x4_t b) { return vqrshrun_high_n_s32(a, b, 9); } @@ -8836,7 +8836,7 @@ int16x8_t test_vqrshrun_high_n_s32(int16x4_t a, int32x4_t b) { // CHECK-NEXT: [[SHUFFLE_I:%.*]] = shufflevector <2 x i32> [[A]], <2 x i32> [[VQRSHRUN_N3]], <4 x i32> <i32 0, i32 1, i32 2, i32 3> // CHECK-NEXT: ret <4 x i32> [[SHUFFLE_I]] // -int32x4_t test_vqrshrun_high_n_s64(int32x2_t a, int64x2_t b) { +uint32x4_t test_vqrshrun_high_n_s64(uint32x2_t a, int64x2_t b) { return vqrshrun_high_n_s64(a, b, 19); } diff --git a/clang/test/CodeGen/X86/prefetchi-error.c b/clang/test/CodeGen/X86/prefetchi-error.c new file mode 100644 index 0000000..31494f7 --- /dev/null +++ b/clang/test/CodeGen/X86/prefetchi-error.c @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 %s -ffreestanding -triple=x86_64-unknown-unknown -target-feature +prefetchi -fsyntax-only -verify + +#include <immintrin.h> + +void test_invalid_prefetchi(void* p) { + __builtin_ia32_prefetchi(p, 1); // expected-error {{argument value 1 is outside the valid range [2, 3]}} +} diff --git a/clang/test/CodeGen/builtin-maximumnum-minimumnum.c b/clang/test/CodeGen/builtin-maximumnum-minimumnum.c new file mode 100644 index 0000000..ea9d2e7 --- /dev/null +++ b/clang/test/CodeGen/builtin-maximumnum-minimumnum.c @@ -0,0 +1,171 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// RUN: %clang_cc1 -x c++ -std=c++20 -disable-llvm-passes -O3 -triple x86_64 %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK + +typedef _Float16 half8 __attribute__((ext_vector_type(8))); +typedef __bf16 bf16x8 __attribute__((ext_vector_type(8))); +typedef float float4 __attribute__((ext_vector_type(4))); +typedef double double2 __attribute__((ext_vector_type(2))); +typedef long double ldouble2 __attribute__((ext_vector_type(2))); + +// CHECK-LABEL: define dso_local noundef <8 x half> @_Z7pfmin16Dv8_DF16_S_( +// CHECK-SAME: <8 x half> noundef [[A:%.*]], <8 x half> noundef [[B:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <8 x half>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <8 x half>, align 16 +// CHECK-NEXT: store <8 x half> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2:![0-9]+]] +// CHECK-NEXT: store <8 x half> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <8 x half>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <8 x half>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <8 x half> @llvm.minimumnum.v8f16(<8 x half> [[TMP0]], <8 x half> [[TMP1]]) +// CHECK-NEXT: ret <8 x half> [[ELT_MINIMUMNUM]] +// +half8 pfmin16(half8 a, half8 b) { + return __builtin_elementwise_minimumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <8 x bfloat> @_Z8pfmin16bDv8_DF16bS_( +// CHECK-SAME: <8 x bfloat> noundef [[A:%.*]], <8 x bfloat> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <8 x bfloat>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <8 x bfloat>, align 16 +// CHECK-NEXT: store <8 x bfloat> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <8 x bfloat> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <8 x bfloat>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <8 x bfloat>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <8 x bfloat> @llvm.minimumnum.v8bf16(<8 x bfloat> [[TMP0]], <8 x bfloat> [[TMP1]]) +// CHECK-NEXT: ret <8 x bfloat> [[ELT_MINIMUMNUM]] +// +bf16x8 pfmin16b(bf16x8 a, bf16x8 b) { + return __builtin_elementwise_minimumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <4 x float> @_Z7pfmin32Dv4_fS_( +// CHECK-SAME: <4 x float> noundef [[A:%.*]], <4 x float> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <4 x float>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <4 x float>, align 16 +// CHECK-NEXT: store <4 x float> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <4 x float> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <4 x float>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <4 x float>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <4 x float> @llvm.minimumnum.v4f32(<4 x float> [[TMP0]], <4 x float> [[TMP1]]) +// CHECK-NEXT: ret <4 x float> [[ELT_MINIMUMNUM]] +// +float4 pfmin32(float4 a, float4 b) { + return __builtin_elementwise_minimumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <2 x double> @_Z7pfmin64Dv2_dS_( +// CHECK-SAME: <2 x double> noundef [[A:%.*]], <2 x double> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <2 x double>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <2 x double>, align 16 +// CHECK-NEXT: store <2 x double> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x double> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <2 x double>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <2 x double>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <2 x double> @llvm.minimumnum.v2f64(<2 x double> [[TMP0]], <2 x double> [[TMP1]]) +// CHECK-NEXT: ret <2 x double> [[ELT_MINIMUMNUM]] +// +double2 pfmin64(double2 a, double2 b) { + return __builtin_elementwise_minimumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <2 x x86_fp80> @_Z7pfmin80Dv2_eS_( +// CHECK-SAME: ptr noundef byval(<2 x x86_fp80>) align 32 [[TMP0:%.*]], ptr noundef byval(<2 x x86_fp80>) align 32 [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <2 x x86_fp80>, align 32 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <2 x x86_fp80>, align 32 +// CHECK-NEXT: [[A:%.*]] = load <2 x x86_fp80>, ptr [[TMP0]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[B:%.*]] = load <2 x x86_fp80>, ptr [[TMP1]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x x86_fp80> [[A]], ptr [[A_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x x86_fp80> [[B]], ptr [[B_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP2:%.*]] = load <2 x x86_fp80>, ptr [[A_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP3:%.*]] = load <2 x x86_fp80>, ptr [[B_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <2 x x86_fp80> @llvm.minimumnum.v2f80(<2 x x86_fp80> [[TMP2]], <2 x x86_fp80> [[TMP3]]) +// CHECK-NEXT: ret <2 x x86_fp80> [[ELT_MINIMUMNUM]] +// +ldouble2 pfmin80(ldouble2 a, ldouble2 b) { + return __builtin_elementwise_minimumnum(a, b); +} + +// CHECK-LABEL: define dso_local noundef <8 x half> @_Z7pfmax16Dv8_DF16_S_( +// CHECK-SAME: <8 x half> noundef [[A:%.*]], <8 x half> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <8 x half>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <8 x half>, align 16 +// CHECK-NEXT: store <8 x half> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <8 x half> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <8 x half>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <8 x half>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MAXIMUMNUM:%.*]] = call <8 x half> @llvm.maximumnum.v8f16(<8 x half> [[TMP0]], <8 x half> [[TMP1]]) +// CHECK-NEXT: ret <8 x half> [[ELT_MAXIMUMNUM]] +// +half8 pfmax16(half8 a, half8 b) { + return __builtin_elementwise_maximumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <8 x bfloat> @_Z8pfmax16bDv8_DF16bS_( +// CHECK-SAME: <8 x bfloat> noundef [[A:%.*]], <8 x bfloat> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <8 x bfloat>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <8 x bfloat>, align 16 +// CHECK-NEXT: store <8 x bfloat> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <8 x bfloat> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <8 x bfloat>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <8 x bfloat>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MAXIMUMNUM:%.*]] = call <8 x bfloat> @llvm.maximumnum.v8bf16(<8 x bfloat> [[TMP0]], <8 x bfloat> [[TMP1]]) +// CHECK-NEXT: ret <8 x bfloat> [[ELT_MAXIMUMNUM]] +// +bf16x8 pfmax16b(bf16x8 a, bf16x8 b) { + return __builtin_elementwise_maximumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <4 x float> @_Z7pfmax32Dv4_fS_( +// CHECK-SAME: <4 x float> noundef [[A:%.*]], <4 x float> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <4 x float>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <4 x float>, align 16 +// CHECK-NEXT: store <4 x float> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <4 x float> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <4 x float>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <4 x float>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MAXIMUMNUM:%.*]] = call <4 x float> @llvm.maximumnum.v4f32(<4 x float> [[TMP0]], <4 x float> [[TMP1]]) +// CHECK-NEXT: ret <4 x float> [[ELT_MAXIMUMNUM]] +// +float4 pfmax32(float4 a, float4 b) { + return __builtin_elementwise_maximumnum(a, b); +} +// CHECK-LABEL: define dso_local noundef <2 x double> @_Z7pfmax64Dv2_dS_( +// CHECK-SAME: <2 x double> noundef [[A:%.*]], <2 x double> noundef [[B:%.*]]) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <2 x double>, align 16 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <2 x double>, align 16 +// CHECK-NEXT: store <2 x double> [[A]], ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x double> [[B]], ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP0:%.*]] = load <2 x double>, ptr [[A_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP1:%.*]] = load <2 x double>, ptr [[B_ADDR]], align 16, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MAXIMUMNUM:%.*]] = call <2 x double> @llvm.maximumnum.v2f64(<2 x double> [[TMP0]], <2 x double> [[TMP1]]) +// CHECK-NEXT: ret <2 x double> [[ELT_MAXIMUMNUM]] +// +double2 pfmax64(double2 a, double2 b) { + return __builtin_elementwise_maximumnum(a, b); +} + +// CHECK-LABEL: define dso_local noundef <2 x x86_fp80> @_Z7pfmax80Dv2_eS_( +// CHECK-SAME: ptr noundef byval(<2 x x86_fp80>) align 32 [[TMP0:%.*]], ptr noundef byval(<2 x x86_fp80>) align 32 [[TMP1:%.*]]) #[[ATTR2]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[A_ADDR:%.*]] = alloca <2 x x86_fp80>, align 32 +// CHECK-NEXT: [[B_ADDR:%.*]] = alloca <2 x x86_fp80>, align 32 +// CHECK-NEXT: [[A:%.*]] = load <2 x x86_fp80>, ptr [[TMP0]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[B:%.*]] = load <2 x x86_fp80>, ptr [[TMP1]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x x86_fp80> [[A]], ptr [[A_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: store <2 x x86_fp80> [[B]], ptr [[B_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP2:%.*]] = load <2 x x86_fp80>, ptr [[A_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[TMP3:%.*]] = load <2 x x86_fp80>, ptr [[B_ADDR]], align 32, !tbaa [[TBAA2]] +// CHECK-NEXT: [[ELT_MINIMUMNUM:%.*]] = call <2 x x86_fp80> @llvm.minimumnum.v2f80(<2 x x86_fp80> [[TMP2]], <2 x x86_fp80> [[TMP3]]) +// CHECK-NEXT: ret <2 x x86_fp80> [[ELT_MINIMUMNUM]] +// +ldouble2 pfmax80(ldouble2 a, ldouble2 b) { + return __builtin_elementwise_minimumnum(a, b); +} + +//. +// CHECK: [[TBAA2]] = !{[[META3:![0-9]+]], [[META3]], i64 0} +// CHECK: [[META3]] = !{!"omnipotent char", [[META4:![0-9]+]], i64 0} +// CHECK: [[META4]] = !{!"Simple C++ TBAA"} +//. diff --git a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp index 3af5a21..1e977dd 100644 --- a/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp +++ b/clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp @@ -105,7 +105,7 @@ void test_mixed() { __builtin_amdgcn_fence( __ATOMIC_SEQ_CST, "workgroup", "local", "local", "global", "local", "local"); } //. -// CHECK: [[META3]] = !{!"amdgpu-as", !"local"} -// CHECK: [[META4]] = !{!"amdgpu-as", !"global"} +// CHECK: [[META3]] = !{!"amdgpu-synchronize-as", !"local"} +// CHECK: [[META4]] = !{!"amdgpu-synchronize-as", !"global"} // CHECK: [[META5]] = !{[[META4]], [[META3]]} //. diff --git a/clang/test/CodeGenCXX/delete.cpp b/clang/test/CodeGenCXX/delete.cpp index d5b0dc6..21b9f8c 100644 --- a/clang/test/CodeGenCXX/delete.cpp +++ b/clang/test/CodeGenCXX/delete.cpp @@ -76,27 +76,45 @@ namespace test1 { ~A(); }; - // CHECK-LABEL: define{{.*}} void @_ZN5test14testEPA10_A20_NS_1AE( - void test(A (*arr)[10][20]) { + // CHECK-LABEL: define{{.*}} void @_ZN5test11fEPA10_A20_NS_1AE( + void f(A (*arr)[10][20]) { delete [] arr; // CHECK: icmp eq ptr [[PTR:%.*]], null // CHECK-NEXT: br i1 - // CHECK: [[BEGIN:%.*]] = getelementptr inbounds [10 x [20 x [[A:%.*]]]], ptr [[PTR]], i32 0, i32 0, i32 0 - // CHECK-NEXT: [[ALLOC:%.*]] = getelementptr inbounds i8, ptr [[BEGIN]], i64 -8 + // CHECK: [[ALLOC:%.*]] = getelementptr inbounds i8, ptr [[PTR]], i64 -8 // CHECK-NEXT: [[COUNT:%.*]] = load i64, ptr [[ALLOC]] - // CHECK: [[END:%.*]] = getelementptr inbounds [[A]], ptr [[BEGIN]], i64 [[COUNT]] - // CHECK-NEXT: [[ISEMPTY:%.*]] = icmp eq ptr [[BEGIN]], [[END]] + // CHECK: [[END:%.*]] = getelementptr inbounds [[A:%.*]], ptr [[PTR]], i64 [[COUNT]] + // CHECK-NEXT: [[ISEMPTY:%.*]] = icmp eq ptr [[PTR]], [[END]] // CHECK-NEXT: br i1 [[ISEMPTY]], // CHECK: [[PAST:%.*]] = phi ptr [ [[END]], {{%.*}} ], [ [[CUR:%.*]], {{%.*}} ] // CHECK-NEXT: [[CUR:%.*]] = getelementptr inbounds [[A]], ptr [[PAST]], i64 -1 // CHECK-NEXT: call void @_ZN5test11AD1Ev(ptr {{[^,]*}} [[CUR]]) - // CHECK-NEXT: [[ISDONE:%.*]] = icmp eq ptr [[CUR]], [[BEGIN]] + // CHECK-NEXT: [[ISDONE:%.*]] = icmp eq ptr [[CUR]], [[PTR]] // CHECK-NEXT: br i1 [[ISDONE]] // CHECK: [[MUL:%.*]] = mul i64 4, [[COUNT]] // CHECK-NEXT: [[SIZE:%.*]] = add i64 [[MUL]], 8 // CHECK-NEXT: call void @_ZdaPvm(ptr noundef [[ALLOC]], i64 noundef [[SIZE]]) } + + // CHECK-LABEL: define{{.*}} void @_ZN5test11gEPA_NS_1AE( + void g(A (*arr)[]) { + delete [] arr; + // CHECK: icmp eq ptr [[PTR:%.*]], null + // CHECK-NEXT: br i1 + + // CHECK: [[ALLOC:%.*]] = getelementptr inbounds i8, ptr [[PTR]], i64 -8 + // CHECK-NEXT: [[COUNT:%.*]] = load i64, ptr [[ALLOC]] + // CHECK: [[END:%.*]] = getelementptr inbounds [[A:%.*]], ptr [[PTR]], i64 [[COUNT]] + // CHECK-NEXT: [[ISEMPTY:%.*]] = icmp eq ptr [[PTR]], [[END]] + // CHECK-NEXT: br i1 [[ISEMPTY]], + // CHECK: [[PAST:%.*]] = phi ptr [ [[END]], {{%.*}} ], [ [[CUR:%.*]], {{%.*}} ] + // CHECK-NEXT: [[CUR:%.*]] = getelementptr inbounds [[A]], ptr [[PAST]], i64 -1 + // CHECK-NEXT: call void @_ZN5test11AD1Ev(ptr {{[^,]*}} [[CUR]]) + // CHECK-NEXT: [[ISDONE:%.*]] = icmp eq ptr [[CUR]], [[PTR]] + // CHECK-NEXT: br i1 [[ISDONE]] + // CHECK: call void @_ZdaPv(ptr noundef [[ALLOC]]) + } } namespace test2 { diff --git a/clang/test/CodeGenCXX/microsoft-abi-eh-async.cpp b/clang/test/CodeGenCXX/microsoft-abi-eh-async.cpp new file mode 100644 index 0000000..b831737 --- /dev/null +++ b/clang/test/CodeGenCXX/microsoft-abi-eh-async.cpp @@ -0,0 +1,209 @@ +// REQUIRES: x86-registered-target + +// RUN: %clang_cl -c --target=x86_64-windows-msvc /EHa -O2 /GS- \ +// RUN: -Xclang=-import-call-optimization \ +// RUN: /clang:-S /clang:-o- -- %s 2>&1 \ +// RUN: | FileCheck %s + +#ifdef __clang__ +#define NO_TAIL __attribute((disable_tail_calls)) +#else +#define NO_TAIL +#endif + +void might_throw(); +void other_func(int x); + +void does_not_throw() noexcept(true); + +extern "C" void __declspec(dllimport) some_dll_import(); + +class HasDtor { + int x; + char foo[40]; + +public: + explicit HasDtor(int x); + ~HasDtor(); +}; + +class BadError { +public: + int errorCode; +}; + +void normal_has_regions() { + // CHECK-LABEL: .def "?normal_has_regions@@YAXXZ" + // CHECK: .seh_endprologue + + // <-- state -1 (none) + { + HasDtor hd{42}; + + // <-- state goes from -1 to 0 + // because state changes, we expect the HasDtor::HasDtor() call to have a NOP + // CHECK: call "??0HasDtor@@QEAA@H@Z" + // CHECK-NEXT: nop + + might_throw(); + // CHECK: call "?might_throw@@YAXXZ" + // CHECK-NEXT: nop + + // <-- state goes from 0 to -1 because we're about to call HasDtor::~HasDtor() + // CHECK: call "??1HasDtor@@QEAA@XZ" + // <-- state -1 + } + + // <-- state -1 + other_func(10); + // CHECK: call "?other_func@@YAXH@Z" + // CHECK-NEXT: nop + // CHECK: .seh_startepilogue + + // <-- state -1 +} + +// This tests a tail call to a destructor. +void case_dtor_arg_empty_body(HasDtor x) +{ + // CHECK-LABEL: .def "?case_dtor_arg_empty_body@@YAXVHasDtor@@@Z" + // CHECK: jmp "??1HasDtor@@QEAA@XZ" +} + +int case_dtor_arg_empty_with_ret(HasDtor x) +{ + // CHECK-LABEL: .def "?case_dtor_arg_empty_with_ret@@YAHVHasDtor@@@Z" + // CHECK: .seh_endprologue + + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK-NOT: nop + + // The call to HasDtor::~HasDtor() should NOT have a NOP because the + // following "mov eax, 100" instruction is in the same EH state. + + return 100; + + // CHECK: mov eax, 100 + // CHECK: .seh_startepilogue + // CHECK: .seh_endepilogue + // CHECK: .seh_endproc +} + +int case_noexcept_dtor(HasDtor x) noexcept(true) +{ + // CHECK: .def "?case_noexcept_dtor@@YAHVHasDtor@@@Z" + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK-NEXT: mov eax, 100 + // CHECK: .seh_startepilogue + return 100; +} + +void case_except_simple_call() NO_TAIL +{ + does_not_throw(); +} +// CHECK-LABEL: .def "?case_except_simple_call@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK-NEXT: call "?does_not_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue +// CHECK: .seh_endproc + +void case_noexcept_simple_call() noexcept(true) NO_TAIL +{ + does_not_throw(); +} +// CHECK-LABEL: .def "?case_noexcept_simple_call@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK-NEXT: call "?does_not_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue +// CHECK: .seh_endepilogue +// CHECK-NEXT: ret +// CHECK-NEXT: .seh_endproc + +// This tests that the destructor is called right before SEH_BeginEpilogue, +// but in a function that has a return value. Loading the return value +// counts as a real instruction, so there is no need for a NOP after the +// dtor call. +int case_dtor_arg_calls_no_throw(HasDtor x) +{ + does_not_throw(); // no NOP expected + return 100; +} +// CHECK-LABEL: .def "?case_dtor_arg_calls_no_throw@@YAHVHasDtor@@@Z" +// CHECK: .seh_endprologue +// CHECK: "?does_not_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK: "??1HasDtor@@QEAA@XZ" +// CHECK-NEXT: mov eax, 100 +// CHECK: .seh_startepilogue +// CHECK: .seh_endproc + +// Check the behavior of CALLs that are at the end of MBBs. If a CALL is within +// a non-null EH state (state -1) and is at the end of an MBB, then we expect +// to find an EH_LABEL after the CALL. This causes us to insert a NOP, which +// is the desired result. +void case_dtor_runs_after_join(int x) { + // CHECK-LABEL: .def "?case_dtor_runs_after_join@@YAXH@Z" + // CHECK: .seh_endprologue + + // <-- EH state -1 + + // ctor call does not need a NOP, because it has real instructions after it + HasDtor hd{42}; + // CHECK: call "??0HasDtor@@QEAA@H@Z" + // CHECK-NEXT: nop + // CHECK: test + + // <-- EH state transition from -1 0 + if (x) { + might_throw(); // <-- NOP expected (at end of BB w/ EH_LABEL) + // CHECK: call "?might_throw@@YAXXZ" + // CHECK-NEXT: nop + } else { + other_func(10); // <-- NOP expected (at end of BB w/ EH_LABEL) + // CHECK: call "?other_func@@YAXH@Z" + // CHECK-NEXT: nop + } + does_not_throw(); + // <-- EH state transition 0 to -1 + // ~HasDtor() runs + + // CHECK: .seh_endproc + + // CHECK: "$ip2state$?case_dtor_runs_after_join@@YAXH@Z": + // CHECK-NEXT: .long [[func_begin:.Lfunc_begin([0-9]+)@IMGREL]] + // CHECK-NEXT: .long -1 + // CHECK-NEXT: .long [[tmp1:.Ltmp([0-9]+)]]@IMGREL + // CHECK-NEXT: .long 0 + // CHECK-NEXT: .long [[tmp2:.Ltmp([0-9]+)]]@IMGREL + // CHECK-NEXT: .long -1 +} + + +// Check the behavior of NOP padding around tail calls. +// We do not expect to insert NOPs around tail calls. +// However, the first call (to other_func()) does get a NOP +// because it comes before .seh_startepilogue. +void case_tail_call_no_eh(bool b) { + // tail call; no NOP padding after JMP + if (b) { + does_not_throw(); + // <-- no NOP here + return; + } + + other_func(20); + // <-- NOP does get inserted here +} +// CHECK-LABEL: .def "?case_tail_call_no_eh@@YAX_N@Z" +// CHECK: test +// CHECK-NEXT: je .LBB +// CHECK: jmp "?does_not_throw@@YAXXZ" +// CHECK-SAME: TAILCALL +// CHECK-NEXT: .LBB +// CHECK-NEXT: mov ecx, 20 +// CHECK-NEXT: jmp "?other_func@@YAXH@Z" +// CHECK-SAME: TAILCALL +// CHECK-NEXT: # -- End function diff --git a/clang/test/CodeGenCXX/microsoft-abi-eh-disabled.cpp b/clang/test/CodeGenCXX/microsoft-abi-eh-disabled.cpp new file mode 100644 index 0000000..744f863 --- /dev/null +++ b/clang/test/CodeGenCXX/microsoft-abi-eh-disabled.cpp @@ -0,0 +1,140 @@ +// REQUIRES: x86-registered-target +// RUN: %clang_cl -c --target=x86_64-windows-msvc -EHs-c- -O2 -GS- \ +// RUN: -Xclang=-import-call-optimization \ +// RUN: -clang:-S -clang:-o- -- %s 2>&1 \ +// RUN: | FileCheck %s + +#ifdef __clang__ +#define NO_TAIL __attribute((disable_tail_calls)) +#else +#define NO_TAIL +#endif + +void might_throw(); +void other_func(int x); + +void does_not_throw() noexcept(true); + +extern "C" void __declspec(dllimport) some_dll_import(); + +class HasDtor { + int x; + char foo[40]; + +public: + explicit HasDtor(int x); + ~HasDtor(); +}; + +void normal_has_regions() { + { + HasDtor hd{42}; + + // because state changes, we expect the HasDtor::HasDtor() call to have a NOP + might_throw(); + } + + other_func(10); +} +// CHECK-LABEL: .def "?normal_has_regions@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK: call "??0HasDtor@@QEAA@H@Z" +// CHECK-NEXT: call "?might_throw@@YAXXZ" +// CHECK-NEXT: mov +// CHECK: call "??1HasDtor@@QEAA@XZ" +// CHECK-NEXT: mov ecx, 10 +// CHECK-NEXT: call "?other_func@@YAXH@Z" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue +// CHECK-NOT: "$ip2state$?normal_has_regions@@YAXXZ" + +// This tests a tail call to a destructor. +void case_dtor_arg_empty_body(HasDtor x) +{ +} +// CHECK-LABEL: .def "?case_dtor_arg_empty_body@@YAXVHasDtor@@@Z" +// CHECK: jmp "??1HasDtor@@QEAA@XZ" + +int case_dtor_arg_empty_with_ret(HasDtor x) +{ + // The call to HasDtor::~HasDtor() should NOT have a NOP because the + // following "mov eax, 100" instruction is in the same EH state. + return 100; +} +// CHECK-LABEL: .def "?case_dtor_arg_empty_with_ret@@YAHVHasDtor@@@Z" +// CHECK: .seh_endprologue +// CHECK: call "??1HasDtor@@QEAA@XZ" +// CHECK-NOT: nop +// CHECK: mov eax, 100 +// CHECK: .seh_startepilogue +// CHECK: .seh_endepilogue +// CHECK: .seh_endproc + +void case_except_simple_call() NO_TAIL +{ + does_not_throw(); +} + +// This tests that the destructor is called right before SEH_BeginEpilogue, +// but in a function that has a return value. +int case_dtor_arg_calls_no_throw(HasDtor x) +{ + does_not_throw(); // no NOP expected + return 100; +} + +// Check the behavior of CALLs that are at the end of MBBs. If a CALL is within +// a non-null EH state (state -1) and is at the end of an MBB, then we expect +// to find an EH_LABEL after the CALL. This causes us to insert a NOP, which +// is the desired result. +void case_dtor_runs_after_join(int x) { + + // ctor call does not need a NOP, because it has real instructions after it + HasDtor hd{42}; + + if (x) { + might_throw(); + } else { + other_func(10); + } + does_not_throw(); + // ~HasDtor() runs +} + +// CHECK-LABEL: .def "?case_dtor_runs_after_join@@YAXH@Z" +// CHECK: .seh_endprologue +// CHECK: call "??0HasDtor@@QEAA@H@Z" +// CHECK-NEXT: test +// CHECK: call "?might_throw@@YAXXZ" +// CHECK-NEXT: jmp +// CHECK: call "?other_func@@YAXH@Z" +// CHECK-NEXT: .LBB +// CHECK: call "?does_not_throw@@YAXXZ" +// CHECK-NEXT: lea +// CHECK-NEXT: call "??1HasDtor@@QEAA@XZ" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue +// CHECK-NOT: "$ip2state$?case_dtor_runs_after_join@@YAXH@Z": + + +// Check the behavior of NOP padding around tail calls. +// We do not expect to insert NOPs around tail calls. +// However, the first call (to other_func()) does get a NOP +// because it comes before .seh_startepilogue. +void case_tail_call_no_eh() { + // ordinary call + other_func(10); + + // tail call; no NOP padding after JMP + does_not_throw(); +} + +// CHECK-LABEL: .def "?case_tail_call_no_eh@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK: call "?other_func@@YAXH@Z" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue +// CHECK: .seh_endepilogue +// CHECK: jmp "?does_not_throw@@YAXXZ" +// CHECK-NOT: nop +// CHECK: .seh_endproc diff --git a/clang/test/CodeGenCXX/microsoft-abi-eh-ip2state.cpp b/clang/test/CodeGenCXX/microsoft-abi-eh-ip2state.cpp new file mode 100644 index 0000000..0b7b406 --- /dev/null +++ b/clang/test/CodeGenCXX/microsoft-abi-eh-ip2state.cpp @@ -0,0 +1,242 @@ +// REQUIRES: x86-registered-target +// RUN: %clang_cl -c --target=x86_64-windows-msvc -O2 -EHsc -GS- \ +// RUN: -Xclang=-import-call-optimization \ +// RUN: -clang:-S -clang:-o- -- %s 2>&1 \ +// RUN: | FileCheck %s + +#ifdef __clang__ +#define NO_TAIL __attribute((disable_tail_calls)) +#else +#define NO_TAIL +#endif + +void might_throw(); +void other_func(int x); + +void does_not_throw() noexcept(true); + +extern "C" void __declspec(dllimport) some_dll_import(); + +class HasDtor { + int x; + char foo[40]; + +public: + explicit HasDtor(int x); + ~HasDtor(); +}; + +class BadError { +public: + int errorCode; +}; + +// Verify that when NOP padding for IP2State is active *and* Import Call +// Optimization is active that we see both forms of NOP padding. +void case_calls_dll_import() NO_TAIL { + some_dll_import(); +} +// CHECK-LABEL: .def "?case_calls_dll_import@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK: .Limpcall{{[0-9]+}}: +// CHECK-NEXT: rex64 +// CHECK-NEXT: call __imp_some_dll_import +// CHECK-NEXT: nop dword ptr {{\[.*\]}} +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue + +void normal_has_regions() { + + // <-- state -1 (none) + { + HasDtor hd{42}; + + // <-- state goes from -1 to 0 + // because state changes, we expect the HasDtor::HasDtor() call to have a NOP + + might_throw(); + + // <-- state goes from 0 to -1 because we're about to call HasDtor::~HasDtor() + // <-- state -1 + } + + // <-- state -1 + other_func(10); + + // <-- state -1 +} +// CHECK-LABEL: .def "?normal_has_regions@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK: call "??0HasDtor@@QEAA@H@Z" +// CHECK-NEXT: nop +// CHECK: call "?might_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK: call "??1HasDtor@@QEAA@XZ" +// CHECK: call "?other_func@@YAXH@Z" +// CHECK-NEXT: nop +// CHECK: .seh_startepilogue + +// This tests a tail call to a destructor. +void case_dtor_arg_empty_body(HasDtor x) +{ +} +// CHECK-LABEL: .def "?case_dtor_arg_empty_body@@YAXVHasDtor@@@Z" +// CHECK: jmp "??1HasDtor@@QEAA@XZ" + +int case_dtor_arg_empty_with_ret(HasDtor x) +{ + // CHECK-LABEL: .def "?case_dtor_arg_empty_with_ret@@YAHVHasDtor@@@Z" + // CHECK: .seh_endprologue + + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK-NOT: nop + + // The call to HasDtor::~HasDtor() should NOT have a NOP because the + // following "mov eax, 100" instruction is in the same EH state. + + return 100; + + // CHECK: mov eax, 100 + // CHECK: .seh_startepilogue + // CHECK: .seh_endepilogue + // CHECK: .seh_endproc +} + +int case_noexcept_dtor(HasDtor x) noexcept(true) +{ + // CHECK: .def "?case_noexcept_dtor@@YAHVHasDtor@@@Z" + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK-NEXT: mov eax, 100 + // CHECK-NEXT: .seh_startepilogue + return 100; +} + +// Simple call of a function that can throw +void case_except_simple_call() NO_TAIL +{ + might_throw(); +} +// CHECK-LABEL: .def "?case_except_simple_call@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK-NEXT: call "?might_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue + +// Simple call of a function that cannot throw, in a noexcept context. +void case_noexcept_simple_call() noexcept(true) NO_TAIL +{ + does_not_throw(); +} +// CHECK-LABEL: .def "?case_noexcept_simple_call@@YAXXZ" +// CHECK: .seh_endprologue +// CHECK-NEXT: call "?does_not_throw@@YAXXZ" +// CHECK-NEXT: nop +// CHECK-NEXT: .seh_startepilogue + + +// This tests that the destructor is called right before SEH_BeginEpilogue, +// but in a function that has a return value. +int case_dtor_arg_calls_no_throw(HasDtor x) +{ + does_not_throw(); // no NOP expected + return 100; +} + +// Check the behavior of CALLs that are at the end of MBBs. If a CALL is within +// a non-null EH state (state -1) and is at the end of an MBB, then we expect +// to find an EH_LABEL after the CALL. This causes us to insert a NOP, which +// is the desired result. +void case_dtor_runs_after_join(int x) { + // CHECK-LABEL: .def "?case_dtor_runs_after_join@@YAXH@Z" + // CHECK: .seh_endprologue + + // <-- EH state -1 + + // ctor call does not need a NOP, because it has real instructions after it + HasDtor hd{42}; + // CHECK: call "??0HasDtor@@QEAA@H@Z" + // CHECK-NEXT: test + + // <-- EH state transition from -1 0 + if (x) { + might_throw(); // <-- NOP expected (at end of BB w/ EH_LABEL) + // CHECK: call "?might_throw@@YAXXZ" + // CHECK-NEXT: nop + } else { + other_func(10); // <-- NOP expected (at end of BB w/ EH_LABEL) + // CHECK: call "?other_func@@YAXH@Z" + // CHECK-NEXT: nop + } + does_not_throw(); + // <-- EH state transition 0 to -1 + // ~HasDtor() runs + + // CHECK: .seh_endproc + + // CHECK: "$ip2state$?case_dtor_runs_after_join@@YAXH@Z": + // CHECK-NEXT: .long [[func_begin:.Lfunc_begin([0-9]+)@IMGREL]] + // CHECK-NEXT: .long -1 + // CHECK-NEXT: .long [[tmp1:.Ltmp([0-9]+)]]@IMGREL + // CHECK-NEXT: .long 0 + // CHECK-NEXT: .long [[tmp2:.Ltmp([0-9]+)]]@IMGREL + // CHECK-NEXT: .long -1 +} + + +// Check the behavior of NOP padding around tail calls. +// We do not expect to insert NOPs around tail calls. +// However, the first call (to other_func()) does get a NOP +// because it comes before .seh_startepilogue. +void case_tail_call_no_eh() { + // CHECK-LABEL: .def "?case_tail_call_no_eh@@YAXXZ" + // CHECK: .seh_endprologue + + // ordinary call + other_func(10); + // CHECK: call "?other_func@@YAXH@Z" + // CHECK-NEXT: nop + + // tail call; no NOP padding after JMP + does_not_throw(); + + // CHECK: .seh_startepilogue + // CHECK: .seh_endepilogue + // CHECK: jmp "?does_not_throw@@YAXXZ" + // CHECK-NOT: nop + // CHECK: .seh_endproc +} + + +// Check the behavior of a try/catch +int case_try_catch() { + // CHECK-LABEL: .def "?case_try_catch@@YAHXZ" + // CHECK: .seh_endprologue + + // Because of the EH_LABELs, the ctor and other_func() get NOPs. + + int result = 0; + try { + // CHECK: call "??0HasDtor@@QEAA@H@Z" + // CHECK-NEXT: nop + HasDtor hd{20}; + + // CHECK: call "?other_func@@YAXH@Z" + // CHECK-NEXT: nop + other_func(10); + + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK: mov + } catch (BadError& e) { + result = 1; + } + return result; + + // CHECK: .seh_endproc + + // CHECK: .def "?dtor$4@?0??case_try_catch@@YAHXZ@4HA" + // CHECK: .seh_endprologue + // CHECK: call "??1HasDtor@@QEAA@XZ" + // CHECK-NEXT: nop + // CHECK: .seh_startepilogue + // CHECK: .seh_endproc +} diff --git a/clang/test/CodeGenHLSL/BasicFeatures/ArrayOutputArguments.hlsl b/clang/test/CodeGenHLSL/BasicFeatures/ArrayOutputArguments.hlsl index bccfaf5..4571649 100644 --- a/clang/test/CodeGenHLSL/BasicFeatures/ArrayOutputArguments.hlsl +++ b/clang/test/CodeGenHLSL/BasicFeatures/ArrayOutputArguments.hlsl @@ -11,7 +11,7 @@ void increment(inout int Arr[2]) { // CHECK-NEXT: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 @{{.*}}, i32 8, i1 false) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[Tmp]], ptr align 4 [[A]], i32 8, i1 false) -// CHECK-NEXT: call void @{{.*}}increment{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]]) +// CHECK-NEXT: call void @{{.*}}increment{{.*}}(ptr noalias noundef align 4 [[Tmp]]) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: [[Idx:%.*]] = getelementptr inbounds [2 x i32], ptr [[A]], i32 0, i32 0 // CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Idx]], align 4 @@ -32,7 +32,7 @@ void fn2(out int Arr[2]) { // CHECK: [[A:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 @{{.*}}, i32 8, i1 false) -// CHECK-NEXT: call void @{{.*}}fn2{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]]) +// CHECK-NEXT: call void @{{.*}}fn2{{.*}}(ptr noalias noundef align 4 [[Tmp]]) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: [[Idx:%.*]] = getelementptr inbounds [2 x i32], ptr [[A]], i32 0, i32 0 // CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Idx]], align 4 @@ -56,7 +56,7 @@ void nestedCall(inout int Arr[2], uint index) { // CHECK-NEXT: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 @{{.*}}, i32 8, i1 false) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[Tmp]], ptr align 4 [[A]], i32 8, i1 false) -// CHECK-NEXT: call void @{{.*}}nestedCall{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]], i32 noundef 0) +// CHECK-NEXT: call void @{{.*}}nestedCall{{.*}}(ptr noalias noundef align 4 [[Tmp]], i32 noundef 0) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: [[Idx:%.*]] = getelementptr inbounds [2 x i32], ptr [[A]], i32 0, i32 1 // CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Idx]], align 4 @@ -70,7 +70,7 @@ export int arrayCall3() { // CHECK-LABEL: outerCall // CHECK: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[Tmp]], ptr align 4 %{{.*}}, i32 8, i1 false) -// CHECK-NEXT: call void {{.*}}increment{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]]) +// CHECK-NEXT: call void {{.*}}increment{{.*}}(ptr noalias noundef align 4 [[Tmp]]) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 {{.*}}, ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: ret void void outerCall(inout int Arr[2]) { @@ -82,7 +82,7 @@ void outerCall(inout int Arr[2]) { // CHECK-NEXT: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 @{{.*}}, i32 8, i1 false) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[Tmp]], ptr align 4 [[A]], i32 8, i1 false) -// CHECK-NEXT: call void @{{.*}}outerCall{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]]) +// CHECK-NEXT: call void @{{.*}}outerCall{{.*}}(ptr noalias noundef align 4 [[Tmp]]) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: [[Idx:%.*]] = getelementptr inbounds [2 x i32], ptr [[A]], i32 0, i32 0 // CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Idx]], align 4 @@ -110,7 +110,7 @@ void outerCall2(inout int Arr[2]) { // CHECK-NEXT: [[Tmp:%.*]] = alloca [2 x i32], align 4 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 @{{.*}}, i32 8, i1 false) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[Tmp]], ptr align 4 [[A]], i32 8, i1 false) -// CHECK-NEXT: call void @{{.*}}outerCall2{{.*}}(ptr noalias noundef byval([2 x i32]) align 4 [[Tmp]]) +// CHECK-NEXT: call void @{{.*}}outerCall2{{.*}}(ptr noalias noundef align 4 [[Tmp]]) // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[A]], ptr align 4 [[Tmp]], i32 8, i1 false) // CHECK-NEXT: [[Idx:%.*]] = getelementptr inbounds [2 x i32], ptr [[A]], i32 0, i32 0 // CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Idx]], align 4 diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl index 75e9710..e96dd66 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl @@ -108,7 +108,7 @@ // GFX1153: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32" // GFX1200: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32" // GFX1201: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32" -// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf16-trans-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+transpose-load-f4f6-insts,+wavefrontsize32 +// GFX1250: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bf16-trans-insts,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+tanh-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+wavefrontsize32 // GFX1103-W64: "target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64" diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl new file mode 100644 index 0000000..f2552d4 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl @@ -0,0 +1,66 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1250 + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef int v4i __attribute__((ext_vector_type(4))); + +// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 1) +// CHECK-GFX1250-NEXT: ret i32 [[TMP0]] +// +int test_amdgcn_global_load_monitor_b32(global int* inptr) +{ + return __builtin_amdgcn_global_load_monitor_b32(inptr, 1); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b64( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.global.load.monitor.b64.v2i32(ptr addrspace(1) [[INPTR:%.*]], i32 10) +// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] +// +v2i test_amdgcn_global_load_monitor_b64(global v2i* inptr) +{ + return __builtin_amdgcn_global_load_monitor_b64(inptr, 10); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b128( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.monitor.b128.v4i32(ptr addrspace(1) [[INPTR:%.*]], i32 22) +// CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] +// +v4i test_amdgcn_global_load_monitor_b128(global v4i* inptr) +{ + return __builtin_amdgcn_global_load_monitor_b128(inptr, 22); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b32( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.flat.load.monitor.b32.i32(ptr [[INPTR:%.*]], i32 27) +// CHECK-GFX1250-NEXT: ret i32 [[TMP0]] +// +int test_amdgcn_flat_load_monitor_b32(int* inptr) +{ + return __builtin_amdgcn_flat_load_monitor_b32(inptr, 27); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b64( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.flat.load.monitor.b64.v2i32(ptr [[INPTR:%.*]], i32 1) +// CHECK-GFX1250-NEXT: ret <2 x i32> [[TMP0]] +// +v2i test_amdgcn_flat_load_monitor_b64(v2i* inptr) +{ + return __builtin_amdgcn_flat_load_monitor_b64(inptr, 1); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_flat_load_monitor_b128( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.flat.load.monitor.b128.v4i32(ptr [[INPTR:%.*]], i32 0) +// CHECK-GFX1250-NEXT: ret <4 x i32> [[TMP0]] +// +v4i test_amdgcn_flat_load_monitor_b128(v4i* inptr) +{ + return __builtin_amdgcn_flat_load_monitor_b128(inptr, 0); +} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl index a21862c..81f39f9 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl @@ -440,6 +440,25 @@ void test_permlane16_swap(global uint2* out, uint old, uint src) { *out = __builtin_amdgcn_permlane16_swap(old, src, false, true); } +// CHECK-LABEL: @test_prefetch( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[FPTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-NEXT: [[GPTR_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[FPTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FPTR_ADDR]] to ptr +// CHECK-NEXT: [[GPTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GPTR_ADDR]] to ptr +// CHECK-NEXT: store ptr [[FPTR:%.*]], ptr [[FPTR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: store ptr addrspace(1) [[GPTR:%.*]], ptr [[GPTR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[FPTR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.amdgcn.flat.prefetch(ptr [[TMP0]], i32 0) +// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr [[GPTR_ADDR_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.amdgcn.global.prefetch(ptr addrspace(1) [[TMP1]], i32 8) +// CHECK-NEXT: ret void +// +void test_prefetch(generic void *fptr, global void *gptr) { + __builtin_amdgcn_flat_prefetch(fptr, 0); + __builtin_amdgcn_global_prefetch(gptr, 8); +} + // CHECK-LABEL: @test_cvt_f32_fp8_e5m3( // CHECK-NEXT: entry: // CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) diff --git a/clang/test/DebugInfo/KeyInstructions/asm.c b/clang/test/DebugInfo/KeyInstructions/asm.c new file mode 100644 index 0000000..2b33016 --- /dev/null +++ b/clang/test/DebugInfo/KeyInstructions/asm.c @@ -0,0 +1,59 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// RUN: %clang_cc1 -triple aarch64 -target-feature +ls64 -O0 -emit-llvm -x c %s -o - -gkey-instructions -debug-info-kind=line-tables-only -gno-column-info | FileCheck %s +// Partially copied from clang/test/CodeGen/AArch64/ls64-inline-asm.c + +// Check the inline asm call and result store are Key and distinct atoms. + +struct foo { unsigned long long x[8]; }; +// CHECK-LABEL: define dso_local void @load( +// CHECK-SAME: ptr noundef [[OUTPUT:%.*]], ptr noundef [[ADDR:%.*]]) #[[ATTR0:[0-9]+]] !dbg [[DBG5:![0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[OUTPUT_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NEXT: store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 8 +// CHECK-NEXT: store ptr [[ADDR]], ptr [[ADDR_ADDR]], align 8 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 8, !dbg [[DBG9:![0-9]+]] +// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ADDR_ADDR]], align 8, !dbg [[DBG9]] +// CHECK-NEXT: [[TMP2:%.*]] = call i512 asm sideeffect "ld64b $0,[$1]", "=r,r,~{memory}"(ptr [[TMP1]]) #[[ATTR1:[0-9]+]], !dbg [[DBG10:![0-9]+]], !srcloc [[META11:![0-9]+]] +// CHECK-NEXT: store i512 [[TMP2]], ptr [[TMP0]], align 8, !dbg [[DBG12:![0-9]+]] +// CHECK-NEXT: ret void, !dbg [[DBG13:![0-9]+]] +// +void load(struct foo *output, void *addr) { + __asm__ volatile ("ld64b %0,[%1]" : "=r" (*output) : "r" (addr) : "memory"); +} + +// CHECK-LABEL: define dso_local void @load2( +// CHECK-SAME: ptr noundef [[OUTPUT:%.*]], ptr noundef [[ADDR:%.*]]) #[[ATTR0]] !dbg [[DBG14:![0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[OUTPUT_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NEXT: store ptr [[OUTPUT]], ptr [[OUTPUT_ADDR]], align 8 +// CHECK-NEXT: store ptr [[ADDR]], ptr [[ADDR_ADDR]], align 8 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[OUTPUT_ADDR]], align 8, !dbg [[DBG15:![0-9]+]] +// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[ADDR_ADDR]], align 8, !dbg [[DBG15]] +// CHECK-NEXT: [[TMP2:%.*]] = call i32 asm sideeffect "ld64b $0,[$1]", "=r,r,~{memory}"(ptr [[TMP1]]) #[[ATTR1]], !dbg [[DBG16:![0-9]+]], !srcloc [[META17:![0-9]+]] +// CHECK-NEXT: store i32 [[TMP2]], ptr [[TMP0]], align 4, !dbg [[DBG18:![0-9]+]] +// CHECK-NEXT: ret void, !dbg [[DBG19:![0-9]+]] +// +void load2(int *output, void *addr) { + __asm__ volatile ("ld64b %0,[%1]" : "=r" (*output) : "r" (addr) : "memory"); +} +//. +// CHECK: [[META0:![0-9]+]] = distinct !DICompileUnit(language: DW_LANG_C11, file: [[META1:![0-9]+]], producer: "{{.*}}clang version {{.*}}", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, splitDebugInlining: false, nameTableKind: None) +// CHECK: [[META1]] = !DIFile(filename: "{{.*}}<stdin>", directory: {{.*}}) +// CHECK: [[DBG5]] = distinct !DISubprogram(name: "load", scope: [[META6:![0-9]+]], file: [[META6]], line: 21, type: [[META7:![0-9]+]], scopeLine: 21, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: [[META0]], keyInstructions: true) +// CHECK: [[META6]] = !DIFile(filename: "{{.*}}asm.c", directory: {{.*}}) +// CHECK: [[META7]] = !DISubroutineType(types: [[META8:![0-9]+]]) +// CHECK: [[META8]] = !{} +// CHECK: [[DBG9]] = !DILocation(line: 22, scope: [[DBG5]]) +// CHECK: [[DBG10]] = !DILocation(line: 22, scope: [[DBG5]], atomGroup: 1, atomRank: 1) +// CHECK: [[META11]] = !{i64 1458} +// CHECK: [[DBG12]] = !DILocation(line: 22, scope: [[DBG5]], atomGroup: 2, atomRank: 1) +// CHECK: [[DBG13]] = !DILocation(line: 23, scope: [[DBG5]], atomGroup: 3, atomRank: 1) +// CHECK: [[DBG14]] = distinct !DISubprogram(name: "load2", scope: [[META6]], file: [[META6]], line: 38, type: [[META7]], scopeLine: 38, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: [[META0]], keyInstructions: true) +// CHECK: [[DBG15]] = !DILocation(line: 39, scope: [[DBG14]]) +// CHECK: [[DBG16]] = !DILocation(line: 39, scope: [[DBG14]], atomGroup: 1, atomRank: 1) +// CHECK: [[META17]] = !{i64 2501} +// CHECK: [[DBG18]] = !DILocation(line: 39, scope: [[DBG14]], atomGroup: 2, atomRank: 1) +// CHECK: [[DBG19]] = !DILocation(line: 40, scope: [[DBG14]], atomGroup: 3, atomRank: 1) +//. diff --git a/clang/test/Driver/amdgpu-hip-system-arch.c b/clang/test/Driver/amdgpu-hip-system-arch.c index 9c27bc0..12e298a 100644 --- a/clang/test/Driver/amdgpu-hip-system-arch.c +++ b/clang/test/Driver/amdgpu-hip-system-arch.c @@ -14,14 +14,14 @@ // RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-new-driver --offload-arch=native --amdgpu-arch-tool=%t/amdgpu_arch_fail -x hip %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR -// NO-OUTPUT-ERROR: error: cannot determine amdgcn architecture{{.*}}; consider passing it via '--offload-arch' +// NO-OUTPUT-ERROR: error: cannot determine hip architecture{{.*}}; consider passing it via '--offload-arch' // case when amdgpu-arch does not return anything with successful execution // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-arch=native --amdgpu-arch-tool=%t/amdgpu_arch_empty -x hip %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=EMPTY-OUTPUT // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-new-driver --offload-arch=native --amdgpu-arch-tool=%t/amdgpu_arch_empty -x hip %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=EMPTY-OUTPUT -// EMPTY-OUTPUT: error: cannot determine amdgcn architecture: No AMD GPU detected in the system; consider passing it via '--offload-arch' +// EMPTY-OUTPUT: error: cannot determine hip architecture: No GPU detected in the system; consider passing it via '--offload-arch' // case when amdgpu-arch returns a gfx906 GPU. // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib --offload-arch=native --amdgpu-arch-tool=%t/amdgpu_arch_gfx906 -x hip %s 2>&1 \ @@ -36,4 +36,4 @@ // RUN: --offload-arch=native --amdgpu-arch-tool=%t/amdgpu_arch_gfx906 \ // RUN: -x hip %s 2>&1 | \ // RUN: FileCheck %s --check-prefix=BAD-TIMEOUT -// BAD-TIMEOUT: clang: error: cannot determine amdgcn architecture: CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected an integer, got 'foo'; consider passing it via '--offload-arch'; environment variable CLANG_TOOLCHAIN_PROGRAM_TIMEOUT specifies the tool timeout (integer secs, <=0 is infinite) +// BAD-TIMEOUT: clang: error: cannot determine hip architecture: CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected an integer, got 'foo'; consider passing it via '--offload-arch'; environment variable CLANG_TOOLCHAIN_PROGRAM_TIMEOUT specifies the tool timeout (integer secs, <=0 is infinite) diff --git a/clang/test/Driver/baremetal.cpp b/clang/test/Driver/baremetal.cpp index 4dc3201..adb59e1 100644 --- a/clang/test/Driver/baremetal.cpp +++ b/clang/test/Driver/baremetal.cpp @@ -257,7 +257,7 @@ // CHECK-RV64-SAME:"{{.*}}.o" // CHECK-RV64-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV64-SAME: "-lc" -// CHECK-RV64-SAME: "-X" "-o" "{{.*}}.tmp.out" +// CHECK-RV64-SAME: "-o" "{{.*}}.tmp.out" // RUN: %clangxx %s -### --target=riscv64-unknown-elf 2>&1 \ // RUN: --sysroot=%S/Inputs/basic_riscv64_tree/riscv64-unknown-elf \ @@ -271,7 +271,7 @@ // CHECK-RV64-DEFAULTCXX-SAME: "-lc++" "-lm" // CHECK-RV64-DEFAULTCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV64-DEFAULTCXX-SAME: "-lc" -// CHECK-RV64-DEFAULTCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV64-DEFAULTCXX-SAME: "-o" "a.out" // RUN: %clangxx %s -### --target=riscv64-unknown-elf 2>&1 \ // RUN: --sysroot=%S/Inputs/basic_riscv64_tree/riscv64-unknown-elf \ @@ -288,7 +288,7 @@ // CHECK-RV64-LIBCXX-SAME: "-lc++" "-lm" // CHECK-RV64-LIBCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV64-LIBCXX-SAME: "-lc" -// CHECK-RV64-LIBCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV64-LIBCXX-SAME: "-o" "a.out" // RUN: %clangxx %s -### 2>&1 --target=riscv64-unknown-elf \ // RUN: --sysroot=%S/Inputs/basic_riscv64_tree/riscv64-unknown-elf \ @@ -305,7 +305,7 @@ // CHECK-RV64-LIBSTDCXX-SAME: "-lstdc++" "-lm" // CHECK-RV64-LIBSTDCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV64-LIBSTDCXX-SAME: "-lc" -// CHECK-RV64-LIBSTDCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV64-LIBSTDCXX-SAME: "-o" "a.out" // RUN: %clang %s -### 2>&1 --target=riscv32-unknown-elf \ // RUN: -L some/directory/user/asked/for \ @@ -325,7 +325,7 @@ // CHECK-RV32-SAME: "{{.*}}.o" // CHECK-RV32-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV32-SAME: "-lc" -// CHECK-RV32-SAME: "-X" "-o" "a.out" +// CHECK-RV32-SAME: "-o" "a.out" // RUN: %clangxx %s -### 2>&1 --target=riscv32-unknown-elf \ // RUN: --sysroot=%S/Inputs/basic_riscv32_tree/riscv32-unknown-elf \ @@ -339,7 +339,7 @@ // CHECK-RV32-DEFAULTCXX-SAME: "-lc++" "-lm" // CHECK-RV32-DEFAULTCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV32-DEFAULTCXX-SAME: "-lc" -// CHECK-RV32-DEFAULTCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV32-DEFAULTCXX-SAME: "-o" "a.out" // RUN: %clangxx %s -### 2>&1 --target=riscv32-unknown-elf \ // RUN: --sysroot=%S/Inputs/basic_riscv32_tree/riscv32-unknown-elf \ @@ -355,7 +355,7 @@ // CHECK-RV32-LIBCXX-SAME: "{{.*}}.o" // CHECK-RV32-LIBCXX-SAME: "-lc++" "-lm" // CHECK-RV32-LIBCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" -// CHECK-RV32-LIBCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV32-LIBCXX-SAME: "-o" "a.out" // RUN: %clangxx %s -### 2>&1 --target=riscv32-unknown-elf \ // RUN: --sysroot=%S/Inputs/basic_riscv32_tree/riscv32-unknown-elf \ @@ -372,7 +372,7 @@ // CHECK-RV32-LIBSTDCXX-SAME: "-lstdc++" "-lm" // CHECK-RV32-LIBSTDCXX-SAME: "{{[^"]*}}libclang_rt.builtins.a" // CHECK-RV32-LIBSTDCXX-SAME: "-lc" -// CHECK-RV32-LIBSTDCXX-SAME: "-X" "-o" "a.out" +// CHECK-RV32-LIBSTDCXX-SAME: "-o" "a.out" // RUN: %clang %s -### 2>&1 --target=riscv64-unknown-elf \ // RUN: -nostdlibinc -nobuiltininc \ diff --git a/clang/test/Driver/cuda-phases.cu b/clang/test/Driver/cuda-phases.cu index 8b91a1d..220a320 100644 --- a/clang/test/Driver/cuda-phases.cu +++ b/clang/test/Driver/cuda-phases.cu @@ -324,8 +324,8 @@ // RUN: -ccc-print-phases --offload-arch=sm_999 -fgpu-rdc -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=INVALID-ARCH %s // INVALID-ARCH: error: unsupported CUDA gpu architecture: sm_999 -// INVALID-ARCH-NEXT: 0: input, "[[INPUT:.+]]", cuda, (host-cuda) -// INVALID-ARCH-NEXT: 1: preprocessor, {0}, cuda-cpp-output, (host-cuda) -// INVALID-ARCH-NEXT: 2: compiler, {1}, ir, (host-cuda) -// INVALID-ARCH-NEXT: 3: backend, {2}, assembler, (host-cuda) -// INVALID-ARCH-NEXT: 4: assembler, {3}, object, (host-cuda) +// INVALID-ARCH: 0: input, "[[INPUT:.+]]", cuda +// INVALID-ARCH-NEXT: 1: preprocessor, {0}, cuda-cpp-output +// INVALID-ARCH-NEXT: 2: compiler, {1}, ir +// INVALID-ARCH-NEXT: 3: backend, {2}, assembler +// INVALID-ARCH-NEXT: 4: assembler, {3}, object diff --git a/clang/test/Driver/fsanitize-ignorelist.c b/clang/test/Driver/fsanitize-ignorelist.c index 7dd666a..d3c8e6c 100644 --- a/clang/test/Driver/fsanitize-ignorelist.c +++ b/clang/test/Driver/fsanitize-ignorelist.c @@ -50,7 +50,7 @@ // Driver properly reports malformed ignorelist files. // RUN: not %clang --target=x86_64-linux-gnu -fsanitize=address -fsanitize-ignorelist=%t.second -fsanitize-ignorelist=%t.bad -fsanitize-ignorelist=%t.good %s -### 2>&1 | FileCheck %s --check-prefix=CHECK-BAD-IGNORELIST -// CHECK-BAD-IGNORELIST: error: malformed sanitizer ignorelist: 'error parsing file '{{.*}}.bad': malformed line 1: 'badline'' +// CHECK-BAD-IGNORELIST: error: failed to parse malformed sanitizer ignorelist: ''{{.*}}.bad': malformed line 1: 'badline' // -fno-sanitize-ignorelist disables all ignorelists specified earlier. // RUN: %clang --target=x86_64-linux-gnu -fsanitize=address -fsanitize-ignorelist=%t.good -fno-sanitize-ignorelist -fsanitize-ignorelist=%t.second %s -### 2>&1 | FileCheck %s --check-prefix=CHECK-ONLY-FIRST-DISABLED --implicit-check-not=-fsanitize-ignorelist= @@ -71,3 +71,10 @@ // CHECK-MISSING-CFI-NO-IGNORELIST-NOT: error: no such file or directory: '{{.*}}cfi_ignorelist.txt' // DELIMITERS: {{^ *"}} + +// Check that a missing file passed to -fsanitize-system-ignorelist triggers a clean error without crashing. +// RUN: not %clang --target=x86_64-linux-gnu -Xclang -fsanitize-system-ignorelist=%t.nonexistent %s -c -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-SYSTEM-IGNORELIST-NOFILE +// CHECK-SYSTEM-IGNORELIST-NOFILE: error: failed to load sanitizer ignorelist file: ''{{.*[\\/]fsanitize-ignorelist\.c\.tmp\.nonexistent}}': {{[Nn]o such file or directory}} +// CHECK-SYSTEM-IGNORELIST-NOFILE-NOT: Stack dump: +// CHECK-SYSTEM-IGNORELIST-NOFILE-NOT: PLEASE submit a bug report +// CHECK-SYSTEM-IGNORELIST-NOFILE-NOT: diagnostic msg: diff --git a/clang/test/Driver/hip-inputs.hip b/clang/test/Driver/hip-inputs.hip index 2d4cc31..a8e25ad 100644 --- a/clang/test/Driver/hip-inputs.hip +++ b/clang/test/Driver/hip-inputs.hip @@ -15,5 +15,5 @@ // RUN: --hip-link %S/Inputs/hip_multiple_inputs/a.cu 2>&1 \ // RUN: | FileCheck -check-prefix=MIX %s -// CHECK-NOT: error: mixed CUDA and HIP compilation is not supported -// MIX: error: mixed CUDA and HIP compilation is not supported +// CHECK-NOT: error: mixed CUDA and HIP offloading compilation is not supported +// MIX: error: mixed CUDA and HIP offloading compilation is not supported diff --git a/clang/test/Driver/hip-invalid-target-id.hip b/clang/test/Driver/hip-invalid-target-id.hip index 555043f..ad942e4 100644 --- a/clang/test/Driver/hip-invalid-target-id.hip +++ b/clang/test/Driver/hip-invalid-target-id.hip @@ -4,7 +4,7 @@ // RUN: --rocm-path=%S/Inputs/rocm \ // RUN: %s 2>&1 | FileCheck -check-prefix=NOPLUS %s -// NOPLUS: error: invalid target ID 'gfx908xnack' +// NOPLUS: error: unsupported HIP gpu architecture: gfx908xnack // RUN: not %clang -### --target=x86_64-linux-gnu \ // RUN: -x hip --offload-arch=gfx900 \ @@ -22,7 +22,7 @@ // RUN: --rocm-path=%S/Inputs/rocm \ // RUN: %s 2>&1 | FileCheck -check-prefix=UNK %s -// UNK: error: invalid target ID 'gfx908:unknown+' +// UNK: error: unsupported HIP gpu architecture: gfx900+xnack // RUN: not %clang -### --target=x86_64-linux-gnu \ // RUN: -x hip --offload-arch=gfx908 \ @@ -31,7 +31,7 @@ // RUN: --rocm-path=%S/Inputs/rocm \ // RUN: %s 2>&1 | FileCheck -check-prefix=MIXED %s -// MIXED: error: invalid target ID 'gfx908:sramecc+:unknown+' +// MIXED: error: unsupported HIP gpu architecture: gfx900+xnack // RUN: not %clang -### --target=x86_64-linux-gnu \ // RUN: -x hip --offload-arch=gfx908 \ @@ -55,7 +55,7 @@ // RUN: --rocm-path=%S/Inputs/rocm \ // RUN: %s 2>&1 | FileCheck -check-prefix=NOCOLON %s -// NOCOLON: error: invalid target ID 'gfx900+xnack' +// NOCOLON: error: unsupported HIP gpu architecture: gfx900+xnack // RUN: not %clang -### --target=x86_64-linux-gnu \ // RUN: -x hip --offload-arch=gfx908 \ diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip index 4fb5571..ba23bc2 100644 --- a/clang/test/Driver/hip-options.hip +++ b/clang/test/Driver/hip-options.hip @@ -115,11 +115,6 @@ // OMP-NOT: "-cc1"{{.*}} "-triple" "amdgcn-amd-amdhsa" {{.*}} "-fopenmp" // OMP: "-cc1"{{.*}} "-triple" "x86_64-unknown-linux-gnu" {{.*}} "-fopenmp" -// RUN: not %clang --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ -// RUN: --offload-arch=gfx906 -fopenmp=libomp -fopenmp-targets=amdgcn %s 2>&1 \ -// RUN: | FileCheck -check-prefix=OMPTGT %s -// OMPTGT: unsupported option '--offload-targets=' for language mode 'HIP' - // Check -Xoffload-linker option is passed to lld. // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpuinc -nogpulib \ diff --git a/clang/test/Driver/hip-phases.hip b/clang/test/Driver/hip-phases.hip index d8a58b7..0ad5d76 100644 --- a/clang/test/Driver/hip-phases.hip +++ b/clang/test/Driver/hip-phases.hip @@ -675,3 +675,25 @@ // DEVICE-ONLY-NEXT: 2: compiler, {1}, ir, (device-hip, gfx90a) // DEVICE-ONLY-NEXT: 3: backend, {2}, ir, (device-hip, gfx90a) // DEVICE-ONLY-NEXT: 4: offload, "device-hip (amdgcn-amd-amdhsa:gfx90a)" {3}, none + +// +// Test the new driver bundling SPIR-V targets. +// +// RUN: %clang -### --target=x86_64-linux-gnu --offload-new-driver -ccc-print-phases \ +// RUN: --offload-device-only --offload-arch=amdgcnspirv,gfx1030 %s 2>&1 \ +// RUN: | FileCheck -check-prefix=SPIRV-ONLY %s +// SPIRV-ONLY: 0: input, "[[INPUT:.+]]", hip, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 1: preprocessor, {0}, hip-cpp-output, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 2: compiler, {1}, ir, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 3: backend, {2}, assembler, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 4: assembler, {3}, object, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 5: linker, {4}, image, (device-hip, gfx1030) +// SPIRV-ONLY-NEXT: 6: offload, "device-hip (amdgcn-amd-amdhsa:gfx1030)" {5}, image +// SPIRV-ONLY-NEXT: 7: input, "[[INPUT]]", hip, (device-hip, amdgcnspirv) +// SPIRV-ONLY-NEXT: 8: preprocessor, {7}, hip-cpp-output, (device-hip, amdgcnspirv) +// SPIRV-ONLY-NEXT: 9: compiler, {8}, ir, (device-hip, amdgcnspirv) +// SPIRV-ONLY-NEXT: 10: backend, {9}, ir, (device-hip, amdgcnspirv) +// SPIRV-ONLY-NEXT: 11: linker, {10}, image, (device-hip, amdgcnspirv) +// SPIRV-ONLY-NEXT: 12: offload, "device-hip (spirv64-amd-amdhsa:amdgcnspirv)" {11}, image +// SPIRV-ONLY-NEXT: 13: linker, {6, 12}, hip-fatbin, (device-hip) +// SPIRV-ONLY-NEXT: 14: offload, "device-hip (amdgcn-amd-amdhsa)" {13}, none diff --git a/clang/test/Driver/invalid-offload-options.cpp b/clang/test/Driver/invalid-offload-options.cpp index 48d5310..6048a3c 100644 --- a/clang/test/Driver/invalid-offload-options.cpp +++ b/clang/test/Driver/invalid-offload-options.cpp @@ -1,29 +1,7 @@ // UNSUPPORTED: system-windows -// RUN: not %clang -### -x hip --target=x86_64-linux-gnu --offload= \ -// RUN: --hip-path=%S/Inputs/hipspv -nogpuinc -nogpulib %s \ -// RUN: 2>&1 | FileCheck --check-prefix=INVALID-TARGET %s // RUN: not %clang -### -x hip --target=x86_64-linux-gnu --offload=foo \ // RUN: --hip-path=%S/Inputs/hipspv -nogpuinc -nogpulib %s \ // RUN: 2>&1 | FileCheck --check-prefix=INVALID-TARGET %s // INVALID-TARGET: error: invalid or unsupported offload target: '{{.*}}' - -// In the future we should be able to specify multiple targets for HIP -// compilation but currently it is not supported. -// -// RUN: not %clang -### -x hip --target=x86_64-linux-gnu --offload=foo,bar \ -// RUN: --hip-path=%S/Inputs/hipspv -nogpuinc -nogpulib %s \ -// RUN: 2>&1 | FileCheck --check-prefix=TOO-MANY-TARGETS %s -// RUN: not %clang -### -x hip --target=x86_64-linux-gnu \ -// RUN: --offload=foo --offload=bar \ -// RUN: --hip-path=%S/Inputs/hipspv -nogpuinc -nogpulib %s \ -// RUN: 2>&1 | FileCheck --check-prefix=TOO-MANY-TARGETS %s - -// TOO-MANY-TARGETS: error: only one offload target is supported - -// RUN: not %clang -### -x hip --target=x86_64-linux-gnu -nogpuinc -nogpulib \ -// RUN: --offload=amdgcn-amd-amdhsa --offload-arch=gfx900 %s \ -// RUN: 2>&1 | FileCheck --check-prefix=OFFLOAD-ARCH-MIX %s - -// OFFLOAD-ARCH-MIX: error: option '--offload-arch' cannot be specified with '--offload' diff --git a/clang/test/Driver/nvptx-cuda-system-arch.c b/clang/test/Driver/nvptx-cuda-system-arch.c index c54eeac..2d4eca8 100644 --- a/clang/test/Driver/nvptx-cuda-system-arch.c +++ b/clang/test/Driver/nvptx-cuda-system-arch.c @@ -16,14 +16,14 @@ // RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-new-driver --offload-arch=native --nvptx-arch-tool=%t/nvptx_arch_fail -x cuda %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR -// NO-OUTPUT-ERROR: error: cannot determine nvptx64 architecture{{.*}}; consider passing it via '--offload-arch' +// NO-OUTPUT-ERROR: error: cannot determine cuda architecture{{.*}}; consider passing it via '--offload-arch' // case when nvptx-arch does not return anything with successful execution // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-arch=native --nvptx-arch-tool=%t/nvptx_arch_empty -x cuda %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=EMPTY-OUTPUT // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-new-driver --offload-arch=native --nvptx-arch-tool=%t/nvptx_arch_empty -x cuda %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=EMPTY-OUTPUT -// EMPTY-OUTPUT: error: cannot determine nvptx64 architecture: No NVIDIA GPU detected in the system; consider passing it via '--offload-arch' +// EMPTY-OUTPUT: error: cannot determine cuda architecture: No GPU detected in the system; consider passing it via '--offload-arch' // case when nvptx-arch does not return anything with successful execution // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpulib --offload-arch=native --nvptx-arch-tool=%t/nvptx_arch_sm_70 -x cuda --cuda-path=%S/Inputs/CUDA_102/usr/local/cuda %s 2>&1 \ @@ -49,4 +49,4 @@ // RUN: --offload-arch=native --nvptx-arch-tool=%t/nvptx_arch_sm_70 \ // RUN: --cuda-path=%S/Inputs/CUDA_102/usr/local/cuda -x cuda %s 2>&1 | \ // RUN: FileCheck %s --check-prefix=BAD-TIMEOUT -// BAD-TIMEOUT: clang: error: cannot determine nvptx64 architecture: CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected an integer, got 'foo'; consider passing it via '--offload-arch'; environment variable CLANG_TOOLCHAIN_PROGRAM_TIMEOUT specifies the tool timeout (integer secs, <=0 is infinite) +// BAD-TIMEOUT: clang: error: cannot determine cuda architecture: CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected an integer, got 'foo'; consider passing it via '--offload-arch'; environment variable CLANG_TOOLCHAIN_PROGRAM_TIMEOUT specifies the tool timeout (integer secs, <=0 is infinite) diff --git a/clang/test/Driver/offload-target.c b/clang/test/Driver/offload-target.c new file mode 100644 index 0000000..23a2cf2 --- /dev/null +++ b/clang/test/Driver/offload-target.c @@ -0,0 +1,22 @@ +// RUN: %clang -### -fsycl --offload-targets=spirv64 -nogpuinc %s -ccc-print-bindings 2>&1 \ +// RUN: | FileCheck %s -check-prefix=SYCL +// SYCL: "spirv64" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[SYCL_BC:.+]]" + +// RUN: %clang -### --offload-targets=amdgcn-amd-amdhsa -nogpulib -nogpuinc -x hip %s -ccc-print-bindings 2>&1 \ +// RUN: | FileCheck %s -check-prefix=HIP +// HIP: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[AMD_OBJ:.+]]" + +// RUN: %clang -### --offload-targets=nvptx64-nvidia-cuda -nogpulib -nogpuinc -x cuda %s -ccc-print-bindings 2>&1 \ +// RUN: | FileCheck %s -check-prefix=CUDA +// CUDA: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[NV_OBJ:.+]]" + +// RUN: %clang -### --offload-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda -fopenmp=libomp \ +// RUN: -Xarch_amdgcn --offload-arch=gfx90a -Xarch_nvptx64 --offload-arch=sm_89 \ +// RUN: -nogpulib -nogpuinc %s -ccc-print-bindings 2>&1 \ +// RUN: | FileCheck %s -check-prefix=OPENMP +// OPENMP: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[AMD_OBJ:.+]]" +// OPENMP: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]"], output: "[[NV_OBJ:.+]]" + +// RUN: %clang -### --offload-targets=spirv64-amd-amdhsa -nogpulib -nogpuinc -x hip %s -ccc-print-bindings 2>&1 \ +// RUN: | FileCheck %s -check-prefix=HIPSPIRV +// HIPSPIRV: "spirv64-amd-amdhsa" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[AMD_OBJ:.+]]" diff --git a/clang/test/Driver/openbsd.c b/clang/test/Driver/openbsd.c index 6639e9d..1f12cfc 100644 --- a/clang/test/Driver/openbsd.c +++ b/clang/test/Driver/openbsd.c @@ -127,9 +127,12 @@ // UNWIND-TABLES: "-funwind-tables=2" // NO-UNWIND-TABLES-NOT: "-funwind-tables=2" -// Check that the -X and --no-relax flags are passed to the linker on riscv64 +// Check that the -X and --no-relax flags are passed to the linker +// RUN: %clang --target=loongarch64-unknown-openbsd -mno-relax -### %s 2>&1 \ +// RUN: | FileCheck --check-prefix=LA64-FLAGS %s // RUN: %clang --target=riscv64-unknown-openbsd -mno-relax -### %s 2>&1 \ // RUN: | FileCheck -check-prefix=RISCV64-FLAGS %s +// LA64-FLAGS: "-X" "--no-relax" // RISCV64-FLAGS: "-X" "--no-relax" // Check passing LTO flags to the linker diff --git a/clang/test/Driver/openmp-offload.c b/clang/test/Driver/openmp-offload.c index 162ff20..64d45f9 100644 --- a/clang/test/Driver/openmp-offload.c +++ b/clang/test/Driver/openmp-offload.c @@ -7,7 +7,7 @@ /// Check whether an invalid OpenMP target is specified: // RUN: not %clang -### -fopenmp=libomp -fopenmp-targets=aaa-bbb-ccc-ddd %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-INVALID-TARGET %s -// CHK-INVALID-TARGET: error: OpenMP target is invalid: 'aaa-bbb-ccc-ddd' +// CHK-INVALID-TARGET: error: invalid or unsupported offload target: 'aaa-bbb-ccc-ddd' /// ########################################################################### @@ -18,15 +18,6 @@ /// ########################################################################### -/// Check error for no -fopenmp option -// RUN: not %clang -### -fopenmp-targets=powerpc64le-ibm-linux-gnu %s 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-NO-FOPENMP %s -// RUN: not %clang -### -fopenmp=libgomp -fopenmp-targets=powerpc64le-ibm-linux-gnu %s 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-NO-FOPENMP %s -// CHK-NO-FOPENMP: error: '-fopenmp-targets' must be used in conjunction with a '-fopenmp' option compatible with offloading; e.g., '-fopenmp=libomp' or '-fopenmp=libiomp5' - -/// ########################################################################### - /// Check warning for duplicate offloading targets. // RUN: %clang -### -ccc-print-phases -fopenmp=libomp -fopenmp-targets=powerpc64le-ibm-linux-gnu,powerpc64le-ibm-linux-gnu %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-DUPLICATES %s diff --git a/clang/test/Driver/openmp-system-arch.c b/clang/test/Driver/openmp-system-arch.c index b18ecf3..167b07a 100644 --- a/clang/test/Driver/openmp-system-arch.c +++ b/clang/test/Driver/openmp-system-arch.c @@ -24,13 +24,7 @@ // RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib -fopenmp=libomp --offload-arch=native \ // RUN: --nvptx-arch-tool=%t/nvptx_arch_empty --amdgpu-arch-tool=%t/amdgpu_arch_empty %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR -// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib -fopenmp=libomp --offload-arch= \ -// RUN: --nvptx-arch-tool=%t/nvptx_arch_fail --amdgpu-arch-tool=%t/amdgpu_arch_fail %s 2>&1 \ -// RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR -// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -nogpulib -fopenmp=libomp --offload-arch= \ -// RUN: --nvptx-arch-tool=%t/nvptx_arch_empty --amdgpu-arch-tool=%t/amdgpu_arch_empty %s 2>&1 \ -// RUN: | FileCheck %s --check-prefix=NO-OUTPUT-ERROR -// NO-OUTPUT-ERROR: error: failed to deduce triple for target architecture 'native'; specify the triple using '-fopenmp-targets' and '-Xopenmp-target' instead +// NO-OUTPUT-ERROR: error: cannot determine openmp architecture // case when amdgpu-arch succeeds. // RUN: %clang -### --target=x86_64-unknown-linux-gnu -nogpulib -fopenmp=libomp --offload-arch=native \ diff --git a/clang/test/Driver/print-multi-selection-flags.c b/clang/test/Driver/print-multi-selection-flags.c index 8cf8f04bb..b1a0a29 100644 --- a/clang/test/Driver/print-multi-selection-flags.c +++ b/clang/test/Driver/print-multi-selection-flags.c @@ -126,3 +126,20 @@ // CHECK-PIE1: -fpie // CHECK-ROPI: -fropi // CHECK-RWPI: -frwpi + +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a -Os | FileCheck --check-prefix=CHECK-OPT-OS %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a -Oz | FileCheck --check-prefix=CHECK-OPT-OZ %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a -O1 | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a -O2 | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=arm-none-eabi -march=armv7a -O3 | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi -Os | FileCheck --check-prefix=CHECK-OPT-OS %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi -Oz | FileCheck --check-prefix=CHECK-OPT-OZ %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi -O1 | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi -O2 | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang -multi-lib-config=%S/Inputs/multilib/empty.yaml -print-multi-flags-experimental --target=aarch64-none-eabi -O3 | FileCheck --check-prefix=CHECK-OPT %s +// CHECK-OPT-OZ: -Oz +// CHECK-OPT-OS: -Os +// CHECK-OPT-NOT: -Oz +// CHECK-OPT-NOT: -Os diff --git a/clang/test/Driver/sparc-target-features.c b/clang/test/Driver/sparc-target-features.c index a839604..bd17da1 100644 --- a/clang/test/Driver/sparc-target-features.c +++ b/clang/test/Driver/sparc-target-features.c @@ -20,6 +20,11 @@ // RUN: %clang --target=sparc -mvis2 %s -### 2>&1 | FileCheck -check-prefix=VIS2 %s // RUN: %clang --target=sparc -mno-vis2 %s -### 2>&1 | FileCheck -check-prefix=NO-VIS2 %s +/// Solaris/SPARC defaults to -mvis2 +// RUN: %clang --target=sparc-sun-solaris2.11 %s -### 2>&1 | FileCheck -check-prefix=VIS2 %s +// RUN: %clang --target=sparc-sun-solaris2.11 -mno-vis2 %s -### 2>&1 | FileCheck -check-prefix=NO-VIS2 %s +// RUN: %clang --target=sparcv9-sun-solaris2.11 %s -### 2>&1 | FileCheck -check-prefix=VIS2 %s +// RUN: %clang --target=sparcv9-sun-solaris2.11 -mno-vis2 %s -### 2>&1 | FileCheck -check-prefix=NO-VIS2 %s // VIS2: "-target-feature" "+vis2" // NO-VIS2: "-target-feature" "-vis2" @@ -34,4 +39,8 @@ // SOFT-QUAD-FLOAT: "-target-feature" "-hard-quad-float" // RUN: %clang --target=sparc -mv8plus %s -### 2>&1 | FileCheck -check-prefix=V8PLUS %s +/// 32-bit Solaris/SPARC defaults to -mv8plus +// RUN: %clang --target=sparc-sun-solaris2.11 %s -### 2>&1 | FileCheck -check-prefix=V8PLUS %s +// RUN: %clang --target=sparc-sun-solaris2.11 -mno-v8plus %s -### 2>&1 | FileCheck -check-prefix=NO-V8PLUS %s // V8PLUS: "-target-feature" "+v8plus" +// NO-V8PLUS-NOT: "-target-feature" "+v8plus" diff --git a/clang/test/Interpreter/fail.cpp b/clang/test/Interpreter/fail.cpp index 633d927..4963df8 100644 --- a/clang/test/Interpreter/fail.cpp +++ b/clang/test/Interpreter/fail.cpp @@ -18,4 +18,11 @@ extern "C" int printf(const char *, ...); int i = 42; auto r1 = printf("i = %d\n", i); // CHECK: i = 42 + +1aap = 42; // expected-error {{invalid digit 'a' in decimal constant}} +1aap = 42; i = 5; // expected-error {{invalid digit 'a' in decimal constant}} + +printf("i = %d\n", i); +// CHECK: i = 42 + %quit diff --git a/clang/test/Interpreter/pretty-print.c b/clang/test/Interpreter/pretty-print.c index 56488a1..e1408c0 100644 --- a/clang/test/Interpreter/pretty-print.c +++ b/clang/test/Interpreter/pretty-print.c @@ -3,7 +3,7 @@ // RUN: cat %s | clang-repl -Xcc -xc | FileCheck %s // RUN: cat %s | clang-repl -Xcc -std=c++11 | FileCheck %s -// UNSUPPORTED: hwasan +// UNSUPPORTED: hwasan, msan char c = 'a'; c diff --git a/clang/test/Interpreter/pretty-print.cpp b/clang/test/Interpreter/pretty-print.cpp index fd79d31..e1036ab 100644 --- a/clang/test/Interpreter/pretty-print.cpp +++ b/clang/test/Interpreter/pretty-print.cpp @@ -1,6 +1,7 @@ // RUN: clang-repl "int i = 10;" 'extern "C" int printf(const char*,...);' \ // RUN: 'auto r1 = printf("i = %d\n", i);' | FileCheck --check-prefix=CHECK-DRIVER %s -// UNSUPPORTED: system-aix +// The test is flaky with asan https://github.com/llvm/llvm-project/pull/148701. +// UNSUPPORTED: system-aix, asan // CHECK-DRIVER: i = 10 // RUN: cat %s | clang-repl -Xcc -std=c++11 -Xcc -fno-delayed-template-parsing | FileCheck %s extern "C" int printf(const char*,...); diff --git a/clang/test/Layout/ms-no-unique-address.cpp b/clang/test/Layout/ms-no-unique-address.cpp index 51cfd9a..fd92056 100644 --- a/clang/test/Layout/ms-no-unique-address.cpp +++ b/clang/test/Layout/ms-no-unique-address.cpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -std=c++2a -fsyntax-only -triple x86_64-windows-msvc -fms-compatibility -fdump-record-layouts %s | FileCheck %s +// RUN: %clang_cc1 -std=c++2a -fsyntax-only -triple x86_64-uefi -fms-compatibility -fdump-record-layouts %s | FileCheck %s namespace Empty { struct A {}; diff --git a/clang/test/Misc/time-passes.c b/clang/test/Misc/time-passes.c index 370f52e..6afbcd4 100644 --- a/clang/test/Misc/time-passes.c +++ b/clang/test/Misc/time-passes.c @@ -19,6 +19,12 @@ // RUN: -ftime-report-json %s -o /dev/null \ // RUN: -mllvm -info-output-file=%t // RUN: cat %t | FileCheck %s --check-prefixes=JSON +// Check that -stats-file-timers only outputs pass time info in the stats file +// and not stderr. +// RUN: %clang_cc1 -emit-obj -O1 \ +// RUN: %s -o /dev/null \ +// RUN: -stats-file=%t -stats-file-timers 2>&1 | count 0 +// RUN: FileCheck %s -input-file=%t -check-prefixes=JSON // TIME: Pass execution timing report // TIME: Total Execution Time: diff --git a/clang/test/OpenMP/copy-gaps-1.cpp b/clang/test/OpenMP/copy-gaps-1.cpp new file mode 100644 index 0000000..3d4fae3 --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-1.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +struct S { + int x; + int y; + int z; + int *p1; + int *p2; +}; + +struct T : public S { + int a; + int b; + int c; +}; + +int main() { + T v; + +#pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c) + { + v.x++; + v.y += 2; + v.z += 3; + v.p1[0] += 4; + v.a += 7; + v.b += 5; + v.c += 6; + } + + return 0; +} + +// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [10 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 32, i64 4, i64 4, i64 4] +// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [10 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [10 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Check for filling of four non-constant size elements here: the whole struct +// size, the (padded) region covering p1 & p2, and the padding at the end of +// struct T. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[P1P2:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[P1P2]], align 8 +// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 2 +// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8 diff --git a/clang/test/OpenMP/copy-gaps-2.cpp b/clang/test/OpenMP/copy-gaps-2.cpp new file mode 100644 index 0000000..5bf603a --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-2.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +struct S { + int x; + int y; + int z; +}; + +struct M : public S { + int mid; +}; + +struct T : public M { + int a; + int b; + int c; +}; + +int main() { + T v; + +#pragma omp target map(tofrom: v, v.y, v.z, v.a) + { + v.y++; + v.z += 2; + v.a += 3; + v.mid += 5; + } + + return 0; +} + +// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 4, i64 4] +// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill four non-constant size elements here: the whole struct size, the region +// covering v.x, the region covering v.mid and the region covering v.b and v.c. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8 +// CHECK-DAG: [[MID:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2 +// CHECK-DAG: store i64 %{{.+}}, ptr [[MID]], align 8 +// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3 +// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8 diff --git a/clang/test/OpenMP/copy-gaps-3.cpp b/clang/test/OpenMP/copy-gaps-3.cpp new file mode 100644 index 0000000..5febb18 --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-3.cpp @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +struct S { + int x; + int y; + int z; +}; + +struct T : public S { + int a; + int b; + int c; +}; + +int main() { + T v; + + // This one should have no gap between v.z & v.a. +#pragma omp target map(tofrom: v, v.y, v.z, v.a) + { + v.y++; + v.z += 2; + v.a += 3; + } + + return 0; +} + +// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [6 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4] +// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [6 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [6 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill three non-constant size elements here: the whole struct size, the region +// covering v.x, and the region covering v.b and v.c. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8 +// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 2 +// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8 diff --git a/clang/test/OpenMP/copy-gaps-4.cpp b/clang/test/OpenMP/copy-gaps-4.cpp new file mode 100644 index 0000000..7060fe3 --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-4.cpp @@ -0,0 +1,48 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +struct S { + int x; + int y; + char z; // Hidden padding after here... +}; + +struct T : public S { + int a; + int b; + int c; +}; + +int main() { + T v; + +#pragma omp target map(tofrom: v, v.y, v.z, v.a) + { + v.y++; + v.z += 2; + v.a += 3; + } + + return 0; +} + +// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4] +// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill four non-constant size elements here: the whole struct size, the region +// covering v.x, the region covering padding after v.z and the region covering +// v.b and v.c. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8 +// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2 +// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8 +// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3 +// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8 diff --git a/clang/test/OpenMP/copy-gaps-5.cpp b/clang/test/OpenMP/copy-gaps-5.cpp new file mode 100644 index 0000000..fae675d --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-5.cpp @@ -0,0 +1,50 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +template<typename C> +struct S { + C x; + C y; + char z; // Hidden padding after here... +}; + +template<typename C> +struct T : public S<C> { + C a; + C b; + C c; +}; + +int main() { + T<int> v; + +#pragma omp target map(tofrom: v, v.y, v.z, v.a) + { + v.y++; + v.z += 2; + v.a += 3; + } + + return 0; +} + +// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4] +// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill four non-constant size elements here: the whole struct size, the region +// covering v.x, the region covering padding after v.z and the region covering +// v.b and v.c. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8 +// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2 +// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8 +// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3 +// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8 diff --git a/clang/test/OpenMP/copy-gaps-6.cpp b/clang/test/OpenMP/copy-gaps-6.cpp new file mode 100644 index 0000000..9c62fde --- /dev/null +++ b/clang/test/OpenMP/copy-gaps-6.cpp @@ -0,0 +1,87 @@ +// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +struct S { + int x; + int *arr; + int y; + int z; +}; + +int main() { + S v; + +#pragma omp target map(tofrom: v, v.x, v.z) + { + v.x++; + v.y += 2; + v.z += 3; + } + +#pragma omp target map(tofrom: v, v.x, v.arr[:1]) + { + v.x++; + v.y += 2; + v.arr[0] += 2; + v.z += 4; + } + +#pragma omp target map(tofrom: v, v.arr[:1]) + { + v.x++; + v.y += 2; + v.arr[0] += 2; + v.z += 4; + } + + return 0; +} + +// CHECK: [[CSTSZ0:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4] +// CHECK: [[CSTTY0:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]] + +// CHECK: [[CSTSZ1:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4] +// CHECK: [[CSTTY1:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]] + +// CHECK: [[CSTSZ2:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 24, i64 4] +// CHECK: [[CSTTY2:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]] + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill two non-constant size elements here: the whole struct size, and the +// region covering v.arr and v.y. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[ARRY:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRY]], align 8 + +// CHECK: call void + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill two non-constant size elements here: the whole struct size, and the +// region covering v.arr, v.y and v.z. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 +// CHECK-DAG: [[ARRYZ:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 1 +// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRYZ]], align 8 + +// CHECK: call void + +// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) +// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 +// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8 +// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [3 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0 + +// Fill one non-constant size element here: the whole struct size. + +// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [3 x i64], ptr [[SIZES]], i32 0, i32 0 +// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8 diff --git a/clang/test/OpenMP/declare_variant_clauses_ast_print.cpp b/clang/test/OpenMP/declare_variant_clauses_ast_print.cpp index c14e19c..e98a23e 100644 --- a/clang/test/OpenMP/declare_variant_clauses_ast_print.cpp +++ b/clang/test/OpenMP/declare_variant_clauses_ast_print.cpp @@ -38,11 +38,11 @@ #ifndef HEADER #define HEADER -void foo_v1(float *AAA, float *BBB, int *I) {return;} -void foo_v2(float *AAA, float *BBB, int *I) {return;} -void foo_v3(float *AAA, float *BBB, int *I) {return;} +void foo_v1(float *AAA, float *BBB, int &CCC, int *I) {return;} +void foo_v2(float *AAA, float *BBB, int &CCC, int *I) {return;} +void foo_v3(float *AAA, float *BBB, int &CCC, int *I) {return;} -//DUMP: FunctionDecl{{.*}} foo 'void (float *, float *, int *)' +//DUMP: FunctionDecl{{.*}} foo 'void (float *, float *, int &, int *)' //DUMP: OMPDeclareVariantAttr{{.*}}device={arch(x86, x86_64)} //DUMP: DeclRefExpr{{.*}}Function{{.*}}foo_v3 //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'I' @@ -54,9 +54,9 @@ void foo_v3(float *AAA, float *BBB, int *I) {return;} //DUMP: DeclRefExpr{{.*}}Function{{.*}}foo_v1 //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'AAA' //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'BBB' -//PRINT: #pragma omp declare variant(foo_v3) match(construct={dispatch}, device={arch(x86, x86_64)}) adjust_args(nothing:I) adjust_args(need_device_ptr:BBB) adjust_args(need_device_addr:AAA) +//PRINT: #pragma omp declare variant(foo_v3) match(construct={dispatch}, device={arch(x86, x86_64)}) adjust_args(nothing:I) adjust_args(need_device_ptr:BBB) adjust_args(need_device_addr:CCC) -//PRINT: #pragma omp declare variant(foo_v2) match(construct={dispatch}, device={arch(ppc)}) adjust_args(need_device_ptr:AAA) adjust_args(need_device_addr:BBB) +//PRINT: #pragma omp declare variant(foo_v2) match(construct={dispatch}, device={arch(ppc)}) adjust_args(need_device_ptr:AAA) adjust_args(need_device_addr:CCC) //PRINT: omp declare variant(foo_v1) match(construct={dispatch}, device={arch(arm)}) adjust_args(need_device_ptr:AAA,BBB) @@ -67,33 +67,33 @@ void foo_v3(float *AAA, float *BBB, int *I) {return;} #pragma omp declare variant(foo_v2) \ match(construct={dispatch}, device={arch(ppc)}), \ adjust_args(need_device_ptr:AAA) \ - adjust_args(need_device_addr:BBB) + adjust_args(need_device_addr:CCC) #pragma omp declare variant(foo_v3) \ adjust_args(need_device_ptr:BBB) adjust_args(nothing:I) \ - adjust_args(need_device_addr:AAA) \ + adjust_args(need_device_addr:CCC) \ match(construct={dispatch}, device={arch(x86,x86_64)}) -void foo(float *AAA, float *BBB, int *I) {return;} +void foo(float *AAA, float *BBB, int &CCC, int *I) {return;} -void Foo_Var(float *AAA, float *BBB, float *CCC) {return;} +void Foo_Var(float *AAA, float *BBB, float *&CCC) {return;} #pragma omp declare variant(Foo_Var) \ match(construct={dispatch}, device={arch(x86_64)}) \ adjust_args(need_device_ptr:AAA) adjust_args(nothing:BBB) \ adjust_args(need_device_addr:CCC) template<typename T> -void Foo(T *AAA, T *BBB, T *CCC) {return;} +void Foo(T *AAA, T *BBB, T *&CCC) {return;} //PRINT: #pragma omp declare variant(Foo_Var) match(construct={dispatch}, device={arch(x86_64)}) adjust_args(nothing:BBB) adjust_args(need_device_ptr:AAA) adjust_args(need_device_addr:CCC) -//DUMP: FunctionDecl{{.*}} Foo 'void (T *, T *, T *)' +//DUMP: FunctionDecl{{.*}} Foo 'void (T *, T *, T *&)' //DUMP: OMPDeclareVariantAttr{{.*}}device={arch(x86_64)} //DUMP: DeclRefExpr{{.*}}Function{{.*}}Foo_Var //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'BBB' //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'AAA' //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'CCC' // -//DUMP: FunctionDecl{{.*}} Foo 'void (float *, float *, float *)' +//DUMP: FunctionDecl{{.*}} Foo 'void (float *, float *, float *&)' //DUMP: OMPDeclareVariantAttr{{.*}}device={arch(x86_64)} //DUMP: DeclRefExpr{{.*}}Function{{.*}}Foo_Var //DUMP: DeclRefExpr{{.*}}ParmVar{{.*}}'BBB' diff --git a/clang/test/OpenMP/declare_variant_clauses_messages.cpp b/clang/test/OpenMP/declare_variant_clauses_messages.cpp index bca9148..916d15f 100644 --- a/clang/test/OpenMP/declare_variant_clauses_messages.cpp +++ b/clang/test/OpenMP/declare_variant_clauses_messages.cpp @@ -91,6 +91,7 @@ void foo_v1(float *AAA, float *BBB, int *I) { return; } void foo_v2(float *AAA, float *BBB, int *I) { return; } void foo_v3(float *AAA, float *BBB, int *I) { return; } void foo_v4(float *AAA, float *BBB, int *I, omp_interop_t IOp) { return; } +void foo_v5(float *AAA, float *BBB, int I) { return; } #if _OPENMP >= 202011 // At least OpenMP 5.1 void vararg_foo(const char *fmt, omp_interop_t it, ...); @@ -129,6 +130,11 @@ void vararg_bar2(const char *fmt) { return; } adjust_args(nothing:J) \ match(construct={dispatch}, device={arch(x86,x86_64)}) +// expected-error@+2 {{expected reference type argument on 'adjust_args' clause with 'need_device_addr' modifier}} +#pragma omp declare variant(foo_v1) \ + adjust_args(need_device_addr:AAA) \ + match(construct={dispatch}, device={arch(x86,x86_64)}) + // expected-error@+2 {{expected reference to one of the parameters of function 'foo'}} #pragma omp declare variant(foo_v3) \ adjust_args(nothing:Other) \ @@ -218,6 +224,12 @@ void vararg_bar2(const char *fmt) { return; } void foo(float *AAA, float *BBB, int *I) { return; } +// expected-error@+2 {{expected reference type argument on 'adjust_args' clause with 'need_device_addr' modifier}} +#pragma omp declare variant(foo_v5) \ + adjust_args(need_device_addr:I) \ + match(construct={dispatch}, device={arch(x86,x86_64)}) +void foo5(float *AAA, float *BBB, int I) { return; } + #endif // NO_INTEROP_T_DEF #ifdef C diff --git a/clang/test/OpenMP/target_map_array_section_no_length_codegen.cpp b/clang/test/OpenMP/target_map_array_section_no_length_codegen.cpp new file mode 100644 index 0000000..43fd509 --- /dev/null +++ b/clang/test/OpenMP/target_map_array_section_no_length_codegen.cpp @@ -0,0 +1,361 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 5 +// RUN: %clang_cc1 -verify -triple i386-unknown-unknown -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s + +// RUN: %clang_cc1 -verify -triple i386-unknown-unknown -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -verify -triple i386-unknown-unknown -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -include-pch %t -emit-llvm %s -o - | FileCheck %s + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER +void array_section_no_length_map_clause(float *d, int index) { + float **f; + + #pragma omp target map(tofrom : d[:]) + { + d[3] += 2; + } + + #pragma omp target map(to : d[2:]) + { + d[3] += 3; + } + + #pragma omp target map(alloc : f[index][:]) + { + f[index][2] += 4; + } + + #pragma omp target map(tofrom : f[index][index+1:]) + { + f[index][index] += 5; + } +} +#endif +// CHECK-LABEL: define dso_local void @_Z34array_section_no_length_map_clausePfi( +// CHECK-SAME: ptr noundef [[D:%.*]], i32 noundef [[INDEX:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: [[INDEX_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[F:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 +// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 4 +// CHECK-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 +// CHECK-NEXT: [[INDEX_CASTED:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS11:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_PTRS12:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS13:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[KERNEL_ARGS14:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 +// CHECK-NEXT: [[INDEX_CASTED17:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS22:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_PTRS23:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS24:%.*]] = alloca [3 x ptr], align 4 +// CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4 +// CHECK-NEXT: [[KERNEL_ARGS25:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8 +// CHECK-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4 +// CHECK-NEXT: store i32 [[INDEX]], ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP2]], i32 0 +// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[TMP1]], ptr [[TMP3]], align 4 +// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP4]], align 4 +// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-NEXT: store ptr null, ptr [[TMP5]], align 4 +// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 +// CHECK-NEXT: store i32 3, ptr [[TMP8]], align 4 +// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 +// CHECK-NEXT: store i32 1, ptr [[TMP9]], align 4 +// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 +// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 4 +// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 +// CHECK-NEXT: store ptr [[TMP7]], ptr [[TMP11]], align 4 +// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 +// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP12]], align 4 +// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 +// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP13]], align 4 +// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 +// CHECK-NEXT: store ptr null, ptr [[TMP14]], align 4 +// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 +// CHECK-NEXT: store ptr null, ptr [[TMP15]], align 4 +// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 +// CHECK-NEXT: store i64 0, ptr [[TMP16]], align 8 +// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 +// CHECK-NEXT: store i64 0, ptr [[TMP17]], align 8 +// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 +// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP18]], align 4 +// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 +// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP19]], align 4 +// CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 +// CHECK-NEXT: store i32 0, ptr [[TMP20]], align 4 +// CHECK-NEXT: [[TMP21:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l14.region_id, ptr [[KERNEL_ARGS]]) +// CHECK-NEXT: [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 0 +// CHECK-NEXT: br i1 [[TMP22]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]] +// CHECK: [[OMP_OFFLOAD_FAILED]]: +// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l14(ptr [[TMP0]]) #[[ATTR2:[0-9]+]] +// CHECK-NEXT: br label %[[OMP_OFFLOAD_CONT]] +// CHECK: [[OMP_OFFLOAD_CONT]]: +// CHECK-NEXT: [[TMP23:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP24:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP25:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw float, ptr [[TMP25]], i32 2 +// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[TMP24]], ptr [[TMP26]], align 4 +// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP27]], align 4 +// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i32 0, i32 0 +// CHECK-NEXT: store ptr null, ptr [[TMP28]], align 4 +// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0 +// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0 +// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0 +// CHECK-NEXT: store i32 3, ptr [[TMP31]], align 4 +// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1 +// CHECK-NEXT: store i32 1, ptr [[TMP32]], align 4 +// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2 +// CHECK-NEXT: store ptr [[TMP29]], ptr [[TMP33]], align 4 +// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3 +// CHECK-NEXT: store ptr [[TMP30]], ptr [[TMP34]], align 4 +// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4 +// CHECK-NEXT: store ptr @.offload_sizes.1, ptr [[TMP35]], align 4 +// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5 +// CHECK-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP36]], align 4 +// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6 +// CHECK-NEXT: store ptr null, ptr [[TMP37]], align 4 +// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7 +// CHECK-NEXT: store ptr null, ptr [[TMP38]], align 4 +// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8 +// CHECK-NEXT: store i64 0, ptr [[TMP39]], align 8 +// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9 +// CHECK-NEXT: store i64 0, ptr [[TMP40]], align 8 +// CHECK-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10 +// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP41]], align 4 +// CHECK-NEXT: [[TMP42:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11 +// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP42]], align 4 +// CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12 +// CHECK-NEXT: store i32 0, ptr [[TMP43]], align 4 +// CHECK-NEXT: [[TMP44:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l19.region_id, ptr [[KERNEL_ARGS5]]) +// CHECK-NEXT: [[TMP45:%.*]] = icmp ne i32 [[TMP44]], 0 +// CHECK-NEXT: br i1 [[TMP45]], label %[[OMP_OFFLOAD_FAILED6:.*]], label %[[OMP_OFFLOAD_CONT7:.*]] +// CHECK: [[OMP_OFFLOAD_FAILED6]]: +// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l19(ptr [[TMP23]]) #[[ATTR2]] +// CHECK-NEXT: br label %[[OMP_OFFLOAD_CONT7]] +// CHECK: [[OMP_OFFLOAD_CONT7]]: +// CHECK-NEXT: [[TMP46:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP47:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: store i32 [[TMP47]], ptr [[INDEX_CASTED]], align 4 +// CHECK-NEXT: [[TMP48:%.*]] = load i32, ptr [[INDEX_CASTED]], align 4 +// CHECK-NEXT: [[TMP49:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP50:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP51:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds ptr, ptr [[TMP50]], i32 [[TMP51]] +// CHECK-NEXT: [[TMP52:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP53:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds ptr, ptr [[TMP52]], i32 [[TMP53]] +// CHECK-NEXT: [[TMP54:%.*]] = load ptr, ptr [[ARRAYIDX9]], align 4 +// CHECK-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds nuw float, ptr [[TMP54]], i32 0 +// CHECK-NEXT: [[TMP55:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[TMP49]], ptr [[TMP55]], align 4 +// CHECK-NEXT: [[TMP56:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[ARRAYIDX8]], ptr [[TMP56]], align 4 +// CHECK-NEXT: [[TMP57:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i32 0, i32 0 +// CHECK-NEXT: store ptr null, ptr [[TMP57]], align 4 +// CHECK-NEXT: [[TMP58:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 1 +// CHECK-NEXT: store ptr [[ARRAYIDX8]], ptr [[TMP58]], align 4 +// CHECK-NEXT: [[TMP59:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 1 +// CHECK-NEXT: store ptr [[ARRAYIDX10]], ptr [[TMP59]], align 4 +// CHECK-NEXT: [[TMP60:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i32 0, i32 1 +// CHECK-NEXT: store ptr null, ptr [[TMP60]], align 4 +// CHECK-NEXT: [[TMP61:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 2 +// CHECK-NEXT: store i32 [[TMP48]], ptr [[TMP61]], align 4 +// CHECK-NEXT: [[TMP62:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 2 +// CHECK-NEXT: store i32 [[TMP48]], ptr [[TMP62]], align 4 +// CHECK-NEXT: [[TMP63:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS13]], i32 0, i32 2 +// CHECK-NEXT: store ptr null, ptr [[TMP63]], align 4 +// CHECK-NEXT: [[TMP64:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS11]], i32 0, i32 0 +// CHECK-NEXT: [[TMP65:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS12]], i32 0, i32 0 +// CHECK-NEXT: [[TMP66:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 0 +// CHECK-NEXT: store i32 3, ptr [[TMP66]], align 4 +// CHECK-NEXT: [[TMP67:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 1 +// CHECK-NEXT: store i32 3, ptr [[TMP67]], align 4 +// CHECK-NEXT: [[TMP68:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 2 +// CHECK-NEXT: store ptr [[TMP64]], ptr [[TMP68]], align 4 +// CHECK-NEXT: [[TMP69:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 3 +// CHECK-NEXT: store ptr [[TMP65]], ptr [[TMP69]], align 4 +// CHECK-NEXT: [[TMP70:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 4 +// CHECK-NEXT: store ptr @.offload_sizes.3, ptr [[TMP70]], align 4 +// CHECK-NEXT: [[TMP71:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 5 +// CHECK-NEXT: store ptr @.offload_maptypes.4, ptr [[TMP71]], align 4 +// CHECK-NEXT: [[TMP72:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 6 +// CHECK-NEXT: store ptr null, ptr [[TMP72]], align 4 +// CHECK-NEXT: [[TMP73:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 7 +// CHECK-NEXT: store ptr null, ptr [[TMP73]], align 4 +// CHECK-NEXT: [[TMP74:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 8 +// CHECK-NEXT: store i64 0, ptr [[TMP74]], align 8 +// CHECK-NEXT: [[TMP75:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 9 +// CHECK-NEXT: store i64 0, ptr [[TMP75]], align 8 +// CHECK-NEXT: [[TMP76:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 10 +// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP76]], align 4 +// CHECK-NEXT: [[TMP77:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 11 +// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP77]], align 4 +// CHECK-NEXT: [[TMP78:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS14]], i32 0, i32 12 +// CHECK-NEXT: store i32 0, ptr [[TMP78]], align 4 +// CHECK-NEXT: [[TMP79:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l24.region_id, ptr [[KERNEL_ARGS14]]) +// CHECK-NEXT: [[TMP80:%.*]] = icmp ne i32 [[TMP79]], 0 +// CHECK-NEXT: br i1 [[TMP80]], label %[[OMP_OFFLOAD_FAILED15:.*]], label %[[OMP_OFFLOAD_CONT16:.*]] +// CHECK: [[OMP_OFFLOAD_FAILED15]]: +// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l24(ptr [[TMP46]], i32 [[TMP48]]) #[[ATTR2]] +// CHECK-NEXT: br label %[[OMP_OFFLOAD_CONT16]] +// CHECK: [[OMP_OFFLOAD_CONT16]]: +// CHECK-NEXT: [[TMP81:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP82:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: store i32 [[TMP82]], ptr [[INDEX_CASTED17]], align 4 +// CHECK-NEXT: [[TMP83:%.*]] = load i32, ptr [[INDEX_CASTED17]], align 4 +// CHECK-NEXT: [[TMP84:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP85:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP86:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX18:%.*]] = getelementptr inbounds ptr, ptr [[TMP85]], i32 [[TMP86]] +// CHECK-NEXT: [[TMP87:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP87]], 1 +// CHECK-NEXT: [[TMP88:%.*]] = load ptr, ptr [[F]], align 4 +// CHECK-NEXT: [[TMP89:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX19:%.*]] = getelementptr inbounds ptr, ptr [[TMP88]], i32 [[TMP89]] +// CHECK-NEXT: [[TMP90:%.*]] = load ptr, ptr [[ARRAYIDX19]], align 4 +// CHECK-NEXT: [[ARRAYIDX20:%.*]] = getelementptr inbounds nuw float, ptr [[TMP90]], i32 [[ADD]] +// CHECK-NEXT: [[TMP91:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ADD21:%.*]] = add nsw i32 [[TMP91]], 1 +// CHECK-NEXT: [[TMP92:%.*]] = mul nuw i32 [[ADD21]], 4 +// CHECK-NEXT: [[TMP93:%.*]] = icmp ugt i32 4, [[TMP92]] +// CHECK-NEXT: [[TMP94:%.*]] = sub nuw i32 4, [[TMP92]] +// CHECK-NEXT: [[TMP95:%.*]] = select i1 [[TMP93]], i32 [[TMP94]], i32 0 +// CHECK-NEXT: [[TMP96:%.*]] = sext i32 [[TMP95]] to i64 +// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[DOTOFFLOAD_SIZES]], ptr align 4 @.offload_sizes.5, i32 24, i1 false) +// CHECK-NEXT: [[TMP97:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS22]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[TMP84]], ptr [[TMP97]], align 4 +// CHECK-NEXT: [[TMP98:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS23]], i32 0, i32 0 +// CHECK-NEXT: store ptr [[ARRAYIDX18]], ptr [[TMP98]], align 4 +// CHECK-NEXT: [[TMP99:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS24]], i32 0, i32 0 +// CHECK-NEXT: store ptr null, ptr [[TMP99]], align 4 +// CHECK-NEXT: [[TMP100:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS22]], i32 0, i32 1 +// CHECK-NEXT: store ptr [[ARRAYIDX18]], ptr [[TMP100]], align 4 +// CHECK-NEXT: [[TMP101:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS23]], i32 0, i32 1 +// CHECK-NEXT: store ptr [[ARRAYIDX20]], ptr [[TMP101]], align 4 +// CHECK-NEXT: [[TMP102:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-NEXT: store i64 [[TMP96]], ptr [[TMP102]], align 4 +// CHECK-NEXT: [[TMP103:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS24]], i32 0, i32 1 +// CHECK-NEXT: store ptr null, ptr [[TMP103]], align 4 +// CHECK-NEXT: [[TMP104:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS22]], i32 0, i32 2 +// CHECK-NEXT: store i32 [[TMP83]], ptr [[TMP104]], align 4 +// CHECK-NEXT: [[TMP105:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS23]], i32 0, i32 2 +// CHECK-NEXT: store i32 [[TMP83]], ptr [[TMP105]], align 4 +// CHECK-NEXT: [[TMP106:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_MAPPERS24]], i32 0, i32 2 +// CHECK-NEXT: store ptr null, ptr [[TMP106]], align 4 +// CHECK-NEXT: [[TMP107:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_BASEPTRS22]], i32 0, i32 0 +// CHECK-NEXT: [[TMP108:%.*]] = getelementptr inbounds [3 x ptr], ptr [[DOTOFFLOAD_PTRS23]], i32 0, i32 0 +// CHECK-NEXT: [[TMP109:%.*]] = getelementptr inbounds [3 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NEXT: [[TMP110:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 0 +// CHECK-NEXT: store i32 3, ptr [[TMP110]], align 4 +// CHECK-NEXT: [[TMP111:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 1 +// CHECK-NEXT: store i32 3, ptr [[TMP111]], align 4 +// CHECK-NEXT: [[TMP112:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 2 +// CHECK-NEXT: store ptr [[TMP107]], ptr [[TMP112]], align 4 +// CHECK-NEXT: [[TMP113:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 3 +// CHECK-NEXT: store ptr [[TMP108]], ptr [[TMP113]], align 4 +// CHECK-NEXT: [[TMP114:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 4 +// CHECK-NEXT: store ptr [[TMP109]], ptr [[TMP114]], align 4 +// CHECK-NEXT: [[TMP115:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 5 +// CHECK-NEXT: store ptr @.offload_maptypes.6, ptr [[TMP115]], align 4 +// CHECK-NEXT: [[TMP116:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 6 +// CHECK-NEXT: store ptr null, ptr [[TMP116]], align 4 +// CHECK-NEXT: [[TMP117:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 7 +// CHECK-NEXT: store ptr null, ptr [[TMP117]], align 4 +// CHECK-NEXT: [[TMP118:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 8 +// CHECK-NEXT: store i64 0, ptr [[TMP118]], align 8 +// CHECK-NEXT: [[TMP119:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 9 +// CHECK-NEXT: store i64 0, ptr [[TMP119]], align 8 +// CHECK-NEXT: [[TMP120:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 10 +// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP120]], align 4 +// CHECK-NEXT: [[TMP121:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 11 +// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP121]], align 4 +// CHECK-NEXT: [[TMP122:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS25]], i32 0, i32 12 +// CHECK-NEXT: store i32 0, ptr [[TMP122]], align 4 +// CHECK-NEXT: [[TMP123:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l29.region_id, ptr [[KERNEL_ARGS25]]) +// CHECK-NEXT: [[TMP124:%.*]] = icmp ne i32 [[TMP123]], 0 +// CHECK-NEXT: br i1 [[TMP124]], label %[[OMP_OFFLOAD_FAILED26:.*]], label %[[OMP_OFFLOAD_CONT27:.*]] +// CHECK: [[OMP_OFFLOAD_FAILED26]]: +// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l29(ptr [[TMP81]], i32 [[TMP83]]) #[[ATTR2]] +// CHECK-NEXT: br label %[[OMP_OFFLOAD_CONT27]] +// CHECK: [[OMP_OFFLOAD_CONT27]]: +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l14( +// CHECK-SAME: ptr noundef [[D:%.*]]) #[[ATTR1:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP0]], i32 3 +// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: [[ADD:%.*]] = fadd float [[TMP1]], 2.000000e+00 +// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l19( +// CHECK-SAME: ptr noundef [[D:%.*]]) #[[ATTR1]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[D_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP0]], i32 3 +// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: [[ADD:%.*]] = fadd float [[TMP1]], 3.000000e+00 +// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l24( +// CHECK-SAME: ptr noundef [[F:%.*]], i32 noundef [[INDEX:%.*]]) #[[ATTR1]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[F_ADDR:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: [[INDEX_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NEXT: store ptr [[F]], ptr [[F_ADDR]], align 4 +// CHECK-NEXT: store i32 [[INDEX]], ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[F_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds ptr, ptr [[TMP0]], i32 [[TMP1]] +// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP2]], i32 2 +// CHECK-NEXT: [[TMP3:%.*]] = load float, ptr [[ARRAYIDX1]], align 4 +// CHECK-NEXT: [[ADD:%.*]] = fadd float [[TMP3]], 4.000000e+00 +// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX1]], align 4 +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z34array_section_no_length_map_clausePfi_l29( +// CHECK-SAME: ptr noundef [[F:%.*]], i32 noundef [[INDEX:%.*]]) #[[ATTR1]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[F_ADDR:%.*]] = alloca ptr, align 4 +// CHECK-NEXT: [[INDEX_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NEXT: store ptr [[F]], ptr [[F_ADDR]], align 4 +// CHECK-NEXT: store i32 [[INDEX]], ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[F_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds ptr, ptr [[TMP0]], i32 [[TMP1]] +// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[ARRAYIDX]], align 4 +// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[INDEX_ADDR]], align 4 +// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP2]], i32 [[TMP3]] +// CHECK-NEXT: [[TMP4:%.*]] = load float, ptr [[ARRAYIDX1]], align 4 +// CHECK-NEXT: [[ADD:%.*]] = fadd float [[TMP4]], 5.000000e+00 +// CHECK-NEXT: store float [[ADD]], ptr [[ARRAYIDX1]], align 4 +// CHECK-NEXT: ret void +// diff --git a/clang/test/OpenMP/target_map_codegen_35.cpp b/clang/test/OpenMP/target_map_codegen_35.cpp index afa0965..c4fc49c 100644 --- a/clang/test/OpenMP/target_map_codegen_35.cpp +++ b/clang/test/OpenMP/target_map_codegen_35.cpp @@ -27,11 +27,11 @@ public: void foo(); }; -// CK35-DAG: [[SIZE_TO:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 0, i64 8] +// CK35-DAG: [[SIZE_TO:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 0, i64 8] // TARGET_PARAM = 0x20 // MEMBER_OF_1 | TO = 0x1000000000001 // MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011 -// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]] +// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]] // CK35-DAG: [[SIZE_FROM:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 8] // TARGET_PARAM = 0x20 // MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012 @@ -86,35 +86,14 @@ void ref_map() { // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint ptr [[B_BEGIN_VOID:%.+]] to i64 // CK35-DAG: [[B_ADDR:%.+]] = getelementptr inbounds nuw %class.S, ptr [[S_ADDR]], i32 0, i32 1 - // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((ptr)(&s+1)-(ptr)(&s.b+1))} to copy the data of remainder of s. + // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b. // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 // CK35-DAG: store ptr [[S_ADDR]], ptr [[BP2]], - // CK35-DAG: store ptr [[B_END:%.+]], ptr [[P2]], - // CK35-DAG: store i64 [[REM_SIZE:%.+]], ptr [[S2]], - - // CK35-DAG: [[B_END]] = getelementptr ptr, ptr [[B_ADDR]], i{{.+}} 1 - - // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) - // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]] - // CK35-DAG: [[B_END_INTPTR]] = ptrtoint ptr [[B_END_VOID:%.+]] to i64 - // CK35-DAG: [[S_END_INTPTR]] = ptrtoint ptr [[S_END_VOID:%.+]] to i64 - // CK35-DAG: [[S_END_VOID]] = getelementptr i8, ptr [[S_LAST:%.+]], i{{.+}} 1 - // CK35-64-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_VOIDPTR:%.+]], i64 15 - // CK35-32-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_VOIDPTR:%.+]], i32 7 - - // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b. - - // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 - // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 - - - // CK35-DAG: store ptr [[S_ADDR]], ptr [[BP3]], - // CK35-DAG: store ptr [[B_ADDR:%.+]], ptr [[P3]], + // CK35-DAG: store ptr [[B_ADDR:%.+]], ptr [[P2]], // CK35-DAG: [[B_ADDR]] = load ptr, ptr [[B_REF:%.+]], align {{[0-9]+}}, !nonnull !{{[0-9]+}}, !align !{{[0-9]+}} // CK35-DAG: [[B_REF]] = getelementptr inbounds nuw %class.S, ptr [[S_ADDR]], i32 0, i32 1 diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp index 4a026584..0ee70be 100644 --- a/clang/test/OpenMP/target_map_messages.cpp +++ b/clang/test/OpenMP/target_map_messages.cpp @@ -122,9 +122,9 @@ struct SA { {} #pragma omp target map(always, tofrom: c,f[1:2]) {} - #pragma omp target map(always, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(always, tofrom: c[:],f) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(always, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(always, tofrom: c,f[:]) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} #pragma omp target map(always) // expected-error {{use of undeclared identifier 'always'}} {} @@ -134,9 +134,9 @@ struct SA { {} #pragma omp target map(self, tofrom: c,f[1:2]) // lt60-error {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} {} - #pragma omp target map(self, tofrom: c[:],f) // lt60-error {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(self, tofrom: c[:],f) // lt60-error {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(self, tofrom: c,f[:]) // lt60-error {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(self, tofrom: c,f[:]) // lt60-error {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} #pragma omp target map(close, tofrom: c,f) {} @@ -144,9 +144,9 @@ struct SA { {} #pragma omp target map(close, tofrom: c,f[1:2]) {} - #pragma omp target map(close, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(close, tofrom: c[:],f) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(close, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(close, tofrom: c,f[:]) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} #pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}} {} @@ -159,11 +159,11 @@ struct SA { // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c,f[1:2]) {} - // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // lt60-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c[:],f) {} - // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // lt60-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c,f[:]) {} @@ -190,14 +190,14 @@ struct SA { {} // ge60-error@+5 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'iterator', 'self}} // ge52-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'iterator'}} - // expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // lt60-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(ompx_hold, tofrom: c[:],f) {} // ge60-error@+5 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'iterator', 'self}} // ge52-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present', 'iterator'}} - // expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // lt60-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(ompx_hold, tofrom: c,f[:]) @@ -448,7 +448,7 @@ void SAclient(int arg) { {} #pragma omp target map(mptr[:1][:2] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} -#pragma omp target map(mptr[:1][:] [0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(mptr[:1][:] [0:2]) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} #pragma omp target map(mptr[:2][:1] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} @@ -517,7 +517,7 @@ void SAclient(int arg) { {} #pragma omp target map(r.S.Ptr [4:5]) {} -#pragma omp target map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(r.S.Ptr[:]) // lt60-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} #pragma omp target map((p + 1)->A) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} diff --git a/clang/test/Preprocessor/arm-acle-6.4.c b/clang/test/Preprocessor/arm-acle-6.4.c index 2c8f486..48deba7 100644 --- a/clang/test/Preprocessor/arm-acle-6.4.c +++ b/clang/test/Preprocessor/arm-acle-6.4.c @@ -188,6 +188,37 @@ // RUN: %clang --target=arm-arm-none-eabi -mcpu=cortex-m33 -x c -E -dM %s -o - | FileCheck %s -check-prefix CHECK-M-DSP // RUN: %clang --target=arm-arm-none-eabi -march=armv8m.main+dsp -x c -E -dM %s -o - | FileCheck %s -check-prefix CHECK-M-DSP +// RUN: %clang -target arm-none-linux-eabi -march=armv8m.base -x c -E -dM %s -o - | FileCheck %s -check-prefix CHECK-V8M-BASE + +// CHECK-V8M-BASE-NOT: __ARM_ARCH_ISA_ARM +// CHECK-V8M-BASE-NOT: __ARM_FEATURE_DSP +// CHECK-V8M-BASE-NOT: __ARM_FEATURE_SIMD32 +// CHECK-V8M-BASE: __ARM_ARCH 8 +// CHECK-V8M-BASE: __ARM_ARCH_ISA_THUMB 1 +// CHECK-V8M-BASE: __ARM_ARCH_PROFILE 'M' +// CHECK-V8M-BASE: __ARM_FEATURE_CLZ 1 +// CHECK-V8M-BASE: __ARM_FEATURE_IDIV 1 +// CHECK-V8M-BASE: __ARM_FEATURE_LDREX 0x7 +// CHECK-V8M-BASE: __ARM_FEATURE_QBIT 1 +// CHECK-V8M-BASE: __ARM_FEATURE_SAT +// CHECK-V8M-BASE-NOT: __ARM_FEATURE_UNALIGNED + +// RUN: %clang -target arm-none-linux-eabi -march=armv8m.main -x c -E -dM %s -o - | FileCheck %s -check-prefix CHECK-V8M-MAIN +// RUN: %clang -target arm-none-linux-eabi -march=armv8.1m.main -x c -E -dM %s -o - | FileCheck %s -check-prefix CHECK-V8M-MAIN + +// CHECK-V8M-MAIN-NOT: __ARM_ARCH_ISA_ARM +// CHECK-V8M-MAIN-NOT: __ARM_FEATURE_DSP +// CHECK-V8M-MAIN-NOT: __ARM_FEATURE_SIMD32 +// CHECK-V8M-MAIN: __ARM_ARCH 8 +// CHECK-V8M-MAIN: __ARM_ARCH_ISA_THUMB 2 +// CHECK-V8M-MAIN: __ARM_ARCH_PROFILE 'M' +// CHECK-V8M-MAIN: __ARM_FEATURE_CLZ 1 +// CHECK-V8M-MAIN: __ARM_FEATURE_IDIV 1 +// CHECK-V8M-MAIN: __ARM_FEATURE_LDREX 0x7 +// CHECK-V8M-MAIN: __ARM_FEATURE_QBIT 1 +// CHECK-V8M-MAIN: __ARM_FEATURE_SAT 1 +// CHECK-V8M-MAIN: __ARM_FEATURE_UNALIGNED 1 + // CHECK-M-DSP: __ARM_FEATURE_DSP 1 // CHECK-M-DSP: __ARM_FEATURE_SIMD32 1 diff --git a/clang/test/Preprocessor/pragma-pushpop-macro-diag.c b/clang/test/Preprocessor/pragma-pushpop-macro-diag.c new file mode 100644 index 0000000..293cb82 --- /dev/null +++ b/clang/test/Preprocessor/pragma-pushpop-macro-diag.c @@ -0,0 +1,4 @@ +// RUN: %clang_cc1 -fms-extensions %s -fsyntax-only -verify + +#pragma push_macro("") // expected-warning {{'#pragma push_macro' expected a non-empty string}} +#pragma pop_macro("") // expected-warning {{'#pragma pop_macro' expected a non-empty string}} diff --git a/clang/test/Preprocessor/pragma-pushpop-macro.c b/clang/test/Preprocessor/pragma-pushpop-macro.c index 0aee074..238e3ed 100644 --- a/clang/test/Preprocessor/pragma-pushpop-macro.c +++ b/clang/test/Preprocessor/pragma-pushpop-macro.c @@ -56,3 +56,6 @@ int P; // CHECK: int pmy2 = 4 // CHECK: int Q; // CHECK: int P; + +#pragma push_macro("") +#pragma pop_macro("") diff --git a/clang/test/Sema/builtins-elementwise-math.c b/clang/test/Sema/builtins-elementwise-math.c index 01057b3..8548d3b 100644 --- a/clang/test/Sema/builtins-elementwise-math.c +++ b/clang/test/Sema/builtins-elementwise-math.c @@ -386,6 +386,96 @@ void test_builtin_elementwise_minimum(int i, short s, float f, double d, float4 // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was '_Complex float')}} } +void test_builtin_elementwise_maximumnum(int i, short s, float f, double d, float4 fv, double4 dv, int3 iv, unsigned3 uv, int *p) { + i = __builtin_elementwise_maximumnum(p, d); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int *')}} + + struct Foo foo = __builtin_elementwise_maximumnum(d, d); + // expected-error@-1 {{initializing 'struct Foo' with an expression of incompatible type 'double'}} + + i = __builtin_elementwise_maximumnum(i); + // expected-error@-1 {{too few arguments to function call, expected 2, have 1}} + + i = __builtin_elementwise_maximumnum(); + // expected-error@-1 {{too few arguments to function call, expected 2, have 0}} + + i = __builtin_elementwise_maximumnum(i, i, i); + // expected-error@-1 {{too many arguments to function call, expected 2, have 3}} + + i = __builtin_elementwise_maximumnum(fv, iv); + // expected-error@-1 {{arguments are of different types ('float4' (vector of 4 'float' values) vs 'int3' (vector of 3 'int' values))}} + + i = __builtin_elementwise_maximumnum(uv, iv); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'unsigned3' (vector of 3 'unsigned int' values))}} + + dv = __builtin_elementwise_maximumnum(fv, dv); + // expected-error@-1 {{arguments are of different types ('float4' (vector of 4 'float' values) vs 'double4' (vector of 4 'double' values))}} + + d = __builtin_elementwise_maximumnum(f, d); + // expected-error@-1 {{arguments are of different types ('float' vs 'double')}} + + fv = __builtin_elementwise_maximumnum(fv, fv); + + i = __builtin_elementwise_maximumnum(iv, iv); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int3' (vector of 3 'int' values))}} + + i = __builtin_elementwise_maximumnum(i, i); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int')}} + + int A[10]; + A = __builtin_elementwise_maximumnum(A, A); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int *')}} + + _Complex float c1, c2; + c1 = __builtin_elementwise_maximumnum(c1, c2); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was '_Complex float')}} +} + +void test_builtin_elementwise_minimumnum(int i, short s, float f, double d, float4 fv, double4 dv, int3 iv, unsigned3 uv, int *p) { + i = __builtin_elementwise_minimumnum(p, d); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int *')}} + + struct Foo foo = __builtin_elementwise_minimumnum(d, d); + // expected-error@-1 {{initializing 'struct Foo' with an expression of incompatible type 'double'}} + + i = __builtin_elementwise_minimumnum(i); + // expected-error@-1 {{too few arguments to function call, expected 2, have 1}} + + i = __builtin_elementwise_minimumnum(); + // expected-error@-1 {{too few arguments to function call, expected 2, have 0}} + + i = __builtin_elementwise_minimumnum(i, i, i); + // expected-error@-1 {{too many arguments to function call, expected 2, have 3}} + + i = __builtin_elementwise_minimumnum(fv, iv); + // expected-error@-1 {{arguments are of different types ('float4' (vector of 4 'float' values) vs 'int3' (vector of 3 'int' values))}} + + i = __builtin_elementwise_minimumnum(uv, iv); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'unsigned3' (vector of 3 'unsigned int' values))}} + + dv = __builtin_elementwise_minimumnum(fv, dv); + // expected-error@-1 {{arguments are of different types ('float4' (vector of 4 'float' values) vs 'double4' (vector of 4 'double' values))}} + + d = __builtin_elementwise_minimumnum(f, d); + // expected-error@-1 {{arguments are of different types ('float' vs 'double')}} + + fv = __builtin_elementwise_minimumnum(fv, fv); + + i = __builtin_elementwise_minimumnum(iv, iv); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int3' (vector of 3 'int' values))}} + + i = __builtin_elementwise_minimumnum(i, i); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int')}} + + int A[10]; + A = __builtin_elementwise_minimumnum(A, A); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was 'int *')}} + + _Complex float c1, c2; + c1 = __builtin_elementwise_minimumnum(c1, c2); + // expected-error@-1 {{1st argument must be a scalar or vector of floating-point types (was '_Complex float')}} +} + void test_builtin_elementwise_bitreverse(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) { struct Foo s = __builtin_elementwise_bitreverse(i); diff --git a/clang/test/Sema/implicit-special-member-deprecated.cpp b/clang/test/Sema/implicit-special-member-deprecated.cpp new file mode 100644 index 0000000..8e23404 --- /dev/null +++ b/clang/test/Sema/implicit-special-member-deprecated.cpp @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -std=c++20 -Wdeprecated-declarations -verify %s + +struct A { + [[deprecated("use something else")]] int x = 42; // expected-note {{marked deprecated here}} +}; + +A makeDefaultA() { return {}; } // ctor is implicit → no warn +A copyA(const A &a) { return a; } // copy-ctor implicit → no warn + +void assignA() { + A a, b; + a = b; // copy-assign implicit → no warn +} + +void useA() { + A a; + (void)a.x; // expected-warning {{is deprecated}} +} + +// Explicitly-defaulted ctor – now silent +struct B { + [[deprecated]] int y; + B() = default; // no warning under new policy +}; diff --git a/clang/test/Sema/unsupported-arm-streaming.cpp b/clang/test/Sema/unsupported-arm-streaming.cpp new file mode 100644 index 0000000..8693dd6 --- /dev/null +++ b/clang/test/Sema/unsupported-arm-streaming.cpp @@ -0,0 +1,3 @@ +// RUN: %clang_cc1 -x c -triple x86_64-pc-linux-gnu -fsyntax-only -verify %s + +void arm_streaming(void) __arm_streaming {} // expected-error {{'__arm_streaming' is not supported on this target}} diff --git a/clang/test/Sema/warn-lifetime-safety-dataflow.cpp b/clang/test/Sema/warn-lifetime-safety-dataflow.cpp index 0e98904ade..2b934ac 100644 --- a/clang/test/Sema/warn-lifetime-safety-dataflow.cpp +++ b/clang/test/Sema/warn-lifetime-safety-dataflow.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -mllvm -debug-only=LifetimeFacts,LifetimeLoanPropagation -Wexperimental-lifetime-safety %s 2>&1 | FileCheck %s +// RUN: %clang_cc1 -fexperimental-lifetime-safety -mllvm -debug-only=LifetimeFacts -Wexperimental-lifetime-safety %s 2>&1 | FileCheck %s // REQUIRES: asserts struct MyObj { @@ -19,10 +19,6 @@ MyObj* return_local_addr() { // CHECK: ReturnOfOrigin (OriginID: [[O_RET_VAL]]) // CHECK: Expire (LoanID: [[L_X]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_ADDR_X]] contains Loan [[L_X]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_X]] -// CHECK-DAG: Origin [[O_RET_VAL]] contains Loan [[L_X]] // Pointer Assignment and Return @@ -47,15 +43,6 @@ MyObj* assign_and_return_local_addr() { // CHECK: ReturnOfOrigin (OriginID: [[O_PTR2_RVAL_2]]) // CHECK: Expire (LoanID: [[L_Y]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_ADDR_Y]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR1]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR2]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR1_RVAL]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR1_RVAL_2]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR2_RVAL]] contains Loan [[L_Y]] -// CHECK-DAG: Origin [[O_PTR2_RVAL_2]] contains Loan [[L_Y]] - // Return of Non-Pointer Type // CHECK-LABEL: Function: return_int_val @@ -65,8 +52,6 @@ int return_int_val() { return x; } // CHECK-NEXT: End of Block -// CHECK: LoanPropagation results: -// CHECK: <empty> // Loan Expiration (Automatic Variable, C++) @@ -79,9 +64,6 @@ void loan_expires_cpp() { // CHECK: AssignOrigin (DestID: [[O_POBJ:[0-9]+]], SrcID: [[O_ADDR_OBJ]]) // CHECK: Expire (LoanID: [[L_OBJ]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_ADDR_OBJ]] contains Loan [[L_OBJ]] -// CHECK-DAG: Origin [[O_POBJ]] contains Loan [[L_OBJ]] // FIXME: No expire for Trivial Destructors @@ -96,10 +78,6 @@ void loan_expires_trivial() { // CHECK-NEXT: End of Block // FIXME: Add check for Expire once trivial destructors are handled for expiration. } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_ADDR_TRIVIAL_OBJ]] contains Loan [[L_TRIVIAL_OBJ]] -// CHECK-DAG: Origin [[O_PTOBJ]] contains Loan [[L_TRIVIAL_OBJ]] - // CHECK-LABEL: Function: conditional void conditional(bool condition) { @@ -119,13 +97,6 @@ void conditional(bool condition) { // CHECK: AssignOrigin (DestID: [[O_P_RVAL:[0-9]+]], SrcID: [[O_P]]) // CHECK: AssignOrigin (DestID: [[O_Q:[0-9]+]], SrcID: [[O_P_RVAL]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_ADDR_A]] contains Loan [[L_A]] -// CHECK-DAG: Origin [[O_ADDR_B]] contains Loan [[L_B]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_A]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_B]] -// CHECK-DAG: Origin [[O_Q]] contains Loan [[L_A]] -// CHECK-DAG: Origin [[O_Q]] contains Loan [[L_B]] // CHECK-LABEL: Function: pointers_in_a_cycle @@ -161,25 +132,6 @@ void pointers_in_a_cycle(bool condition) { // CHECK: AssignOrigin (DestID: [[O_P3]], SrcID: [[O_TEMP_RVAL]]) } } -// At the end of the analysis, the origins for the pointers involved in the cycle -// (p1, p2, p3, temp) should all contain the loans from v1, v2, and v3 at the fixed point. -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_P1]] contains Loan [[L_V1]] -// CHECK-DAG: Origin [[O_P1]] contains Loan [[L_V2]] -// CHECK-DAG: Origin [[O_P1]] contains Loan [[L_V3]] -// CHECK-DAG: Origin [[O_P2]] contains Loan [[L_V1]] -// CHECK-DAG: Origin [[O_P2]] contains Loan [[L_V2]] -// CHECK-DAG: Origin [[O_P2]] contains Loan [[L_V3]] -// CHECK-DAG: Origin [[O_P3]] contains Loan [[L_V1]] -// CHECK-DAG: Origin [[O_P3]] contains Loan [[L_V2]] -// CHECK-DAG: Origin [[O_P3]] contains Loan [[L_V3]] -// CHECK-DAG: Origin [[O_TEMP]] contains Loan [[L_V1]] -// CHECK-DAG: Origin [[O_TEMP]] contains Loan [[L_V2]] -// CHECK-DAG: Origin [[O_TEMP]] contains Loan [[L_V3]] -// CHECK-DAG: Origin [[O_ADDR_V1]] contains Loan [[L_V1]] -// CHECK-DAG: Origin [[O_ADDR_V2]] contains Loan [[L_V2]] -// CHECK-DAG: Origin [[O_ADDR_V3]] contains Loan [[L_V3]] - // CHECK-LABEL: Function: overwrite_origin void overwrite_origin() { @@ -195,10 +147,6 @@ void overwrite_origin() { // CHECK: Expire (LoanID: [[L_S2]]) // CHECK: Expire (LoanID: [[L_S1]]) } -// CHECK: LoanPropagation results: -// CHECK: Origin [[O_P]] contains Loan [[L_S2]] -// CHECK-NOT: Origin [[O_P]] contains Loan [[L_S1]] - // CHECK-LABEL: Function: reassign_to_null void reassign_to_null() { @@ -213,8 +161,6 @@ void reassign_to_null() { } // FIXME: Have a better representation for nullptr than just an empty origin. // It should be a separate loan and origin kind. -// CHECK: LoanPropagation results: -// CHECK: Origin [[O_P]] contains no loans // CHECK-LABEL: Function: reassign_in_if @@ -235,11 +181,6 @@ void reassign_in_if(bool condition) { // CHECK: Expire (LoanID: [[L_S2]]) // CHECK: Expire (LoanID: [[L_S1]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S2]] -// CHECK-DAG: Origin [[O_ADDR_S1]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_ADDR_S2]] contains Loan [[L_S2]] // CHECK-LABEL: Function: assign_in_switch @@ -276,14 +217,6 @@ void assign_in_switch(int mode) { // CHECK-DAG: Expire (LoanID: [[L_S2]]) // CHECK-DAG: Expire (LoanID: [[L_S1]]) } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S2]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S3]] -// CHECK-DAG: Origin [[O_ADDR_S1]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_ADDR_S2]] contains Loan [[L_S2]] -// CHECK-DAG: Origin [[O_ADDR_S3]] contains Loan [[L_S3]] - // CHECK-LABEL: Function: loan_in_loop void loan_in_loop(bool condition) { @@ -299,10 +232,6 @@ void loan_in_loop(bool condition) { // CHECK: Expire (LoanID: [[L_INNER]]) } } -// CHECK: LoanPropagation results: -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_INNER]] -// CHECK-DAG: Origin [[O_ADDR_INNER]] contains Loan [[L_INNER]] - // CHECK-LABEL: Function: loop_with_break void loop_with_break(int count) { @@ -326,13 +255,6 @@ void loop_with_break(int count) { // CHECK: Expire (LoanID: [[L_S1]]) } -// CHECK-LABEL: LoanPropagation results: -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_S2]] -// CHECK-DAG: Origin [[O_ADDR_S1]] contains Loan [[L_S1]] -// CHECK-DAG: Origin [[O_ADDR_S2]] contains Loan [[L_S2]] - - // CHECK-LABEL: Function: nested_scopes void nested_scopes() { MyObj* p = nullptr; @@ -355,13 +277,6 @@ void nested_scopes() { // CHECK: Expire (LoanID: [[L_OUTER]]) } -// CHECK-LABEL: LoanPropagation results: -// CHECK-DAG: Origin [[O_P]] contains Loan [[L_INNER]] -// CHECK-DAG: Origin [[O_ADDR_INNER]] contains Loan [[L_INNER]] -// CHECK-DAG: Origin [[O_ADDR_OUTER]] contains Loan [[L_OUTER]] -// CHECK-NOT: Origin [[O_P]] contains Loan [[L_OUTER]] - - // CHECK-LABEL: Function: pointer_indirection void pointer_indirection() { int a; diff --git a/clang/test/SemaCXX/attr-target-clones-riscv.cpp b/clang/test/SemaCXX/attr-target-clones-riscv.cpp index 102bb4b..7648284 100644 --- a/clang/test/SemaCXX/attr-target-clones-riscv.cpp +++ b/clang/test/SemaCXX/attr-target-clones-riscv.cpp @@ -9,6 +9,9 @@ void __attribute__((target_clones("default", "mtune=sifive-u74"))) mtune() {} // expected-warning@+1 {{version list contains duplicate entries}} void __attribute__((target_clones("default", "arch=+c", "arch=+c"))) dupVersion() {} +// expected-warning@+1 {{version list contains duplicate entries}} +void __attribute__((target_clones(" default", "default "))) dupDefault() {} + // expected-warning@+1 {{unsupported '' in the 'target_clones' attribute string; 'target_clones' attribute ignored}} void __attribute__((target_clones("default", ""))) emptyVersion() {} diff --git a/clang/test/SemaCXX/constant-expression-p2280r4.cpp b/clang/test/SemaCXX/constant-expression-p2280r4.cpp index 16f5f82..312a778 100644 --- a/clang/test/SemaCXX/constant-expression-p2280r4.cpp +++ b/clang/test/SemaCXX/constant-expression-p2280r4.cpp @@ -383,3 +383,17 @@ namespace enable_if_2 { } } } + +namespace GH150015 { + extern int (& c)[8]; // interpreter-note {{declared here}} + constexpr int x = c <= c+8; // interpreter-error {{constexpr variable 'x' must be initialized by a constant expression}} \ + // interpreter-note {{initializer of 'c' is unknown}} + + struct X {}; + struct Y {}; + struct Z : X, Y {}; + extern Z &z; // interpreter-note{{declared here}} + constexpr int bases = (void*)(X*)&z <= (Y*)&z; // expected-error {{constexpr variable 'bases' must be initialized by a constant expression}} \ + // nointerpreter-note {{comparison of addresses of subobjects of different base classes has unspecified value}} \ + // interpreter-note {{initializer of 'z' is unknown}} +} diff --git a/clang/test/SemaCXX/cxx2a-ms-no-unique-address.cpp b/clang/test/SemaCXX/cxx2a-ms-no-unique-address.cpp index 726cfa3..a064401 100644 --- a/clang/test/SemaCXX/cxx2a-ms-no-unique-address.cpp +++ b/clang/test/SemaCXX/cxx2a-ms-no-unique-address.cpp @@ -1,5 +1,6 @@ // RUN: %clang_cc1 -std=c++2a %s -verify=unsupported -triple x86_64-linux-gnu // RUN: %clang_cc1 -std=c++2a %s -verify -triple x86_64-windows -fms-compatibility +// RUN: %clang_cc1 -std=c++2a %s -verify -triple x86_64-uefi -fms-compatibility [[msvc::no_unique_address]] int a; // expected-error {{only applies to non-bit-field non-static data members}} unsupported-warning {{unknown}} [[msvc::no_unique_address]] void f(); // expected-error {{only applies to non-bit-field non-static data members}} unsupported-warning {{unknown}} diff --git a/clang/test/SemaCXX/wreturn-always-throws.cpp b/clang/test/SemaCXX/wreturn-always-throws.cpp index addcadd..df7689f 100644 --- a/clang/test/SemaCXX/wreturn-always-throws.cpp +++ b/clang/test/SemaCXX/wreturn-always-throws.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsyntax-only -fcxx-exceptions -fexceptions -Wreturn-type -verify %s +// RUN: %clang_cc1 -fsyntax-only -fcxx-exceptions -fexceptions -Wreturn-type -Winvalid-noreturn -verify %s // expected-no-diagnostics namespace std { @@ -44,3 +44,22 @@ void testTemplates() { throwErrorTemplate("ERROR"); (void)ensureZeroTemplate(42); } + +// Ensure that explicit specialization of a member function does not inherit +// the warning from the primary template. + +template<typename T> +struct S { + void f(); + void g(); +}; + +template<typename T> +void S<T>::f() { throw 0; } +template<> +void S<int>::f() {} + +template<typename T> +void S<T>::g() {} +template<> +void S<int>::g() { throw 0; } diff --git a/clang/test/SemaOpenACC/atomic-construct.cpp b/clang/test/SemaOpenACC/atomic-construct.cpp index eba2559..10c85c2 100644 --- a/clang/test/SemaOpenACC/atomic-construct.cpp +++ b/clang/test/SemaOpenACC/atomic-construct.cpp @@ -1098,7 +1098,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS--; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used in unary expression('LHS') from the first statement}} LHS = RHS; } @@ -1128,7 +1128,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS *= 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of compound assignment('LHS') from the first statement}} LHS = RHS; } #pragma acc atomic capture @@ -1157,7 +1157,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS = LHS * 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1186,7 +1186,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS = LHS | 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1255,7 +1255,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression on left hand side of assignment('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS = 1; } @@ -1294,7 +1294,7 @@ void AtomicCaptureTemplateCompound(T LHS, T RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression in unary expression('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS++; } } @@ -1352,7 +1352,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS--; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used in unary expression('LHS') from the first statement}} LHS = RHS; } @@ -1384,7 +1384,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS *= 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of compound assignment('LHS') from the first statement}} LHS = RHS; } #pragma acc atomic capture @@ -1415,7 +1415,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS = LHS * 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1446,7 +1446,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS = LHS | 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1523,7 +1523,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression on left hand side of assignment('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS = 1; } @@ -1570,7 +1570,7 @@ void AtomicCaptureTemplateCompound2(T LHS, T RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression in unary expression('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS++; } } @@ -1629,7 +1629,7 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS--; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used in unary expression('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used in unary expression('LHS') from the first statement}} LHS = RHS; } @@ -1666,7 +1666,7 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS *= 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of compound assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of compound assignment('LHS') from the first statement}} LHS = RHS; } #pragma acc atomic capture @@ -1702,7 +1702,7 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS = LHS * 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1738,7 +1738,7 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS = LHS | 1; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on right hand side of assignment('RHS') must match variable used on left hand side of assignment('LHS') from the first statement}} + // expected-note@+1{{sub-expression on right hand side of assignment('RHS') must match sub-expression used on left hand side of assignment('LHS') from the first statement}} RHS = RHS; } #pragma acc atomic capture @@ -1815,7 +1815,7 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable on left hand side of assignment('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression on left hand side of assignment('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS = 1; } @@ -1861,7 +1861,23 @@ void AtomicCaptureCompound(int LHS, int RHS) { { LHS = RHS; // expected-error@-3{{statement associated with OpenACC 'atomic capture' directive is invalid}} - // expected-note@+1{{variable in unary expression('LHS') must match variable used on right hand side of assignment('RHS') from the first statement}} + // expected-note@+1{{sub-expression in unary expression('LHS') must match sub-expression used on right hand side of assignment('RHS') from the first statement}} LHS++; } + + // Example from UDel test suite, which wasn't working because of irrelevant + // parens, make sure we work with these. This should not diagnose. + typedef double real_t; + int * distribution; + real_t *a; + real_t *b; + int *c; + for (int x = 0; x < 5; ++x) { +#pragma acc atomic capture + { + c[x] = distribution[(int) (a[x]*b[x]/10)]; + (distribution[(int)(a[x]*b[x]/10)])--; + } + } + } diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl index 9711b3b..3247380 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl @@ -1,6 +1,7 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1250 -verify -S -o - %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-- -target-cpu gfx1250 -verify -S -o - %s +typedef int v2i __attribute__((ext_vector_type(2))); typedef int v4i __attribute__((ext_vector_type(4))); typedef int v8i __attribute__((ext_vector_type(8))); @@ -28,6 +29,17 @@ void test__builtin_amdgcn_cvt_f16_bf8(int a, int b) { __builtin_amdgcn_cvt_f16_bf8(a, b); // expected-error {{'__builtin_amdgcn_cvt_f16_bf8' must be a constant integer}} } +void test_amdgcn_load_monitor(global int* b32gaddr, global v2i* b64gaddr, global v4i* b128gaddr, int *b32faddr, v2i* b64faddr, v4i *b128faddr, + global int* b32out, global v2i* b64out, global v4i* b128out, int cpol) +{ + *b32out = __builtin_amdgcn_global_load_monitor_b32(b32gaddr, cpol); // expected-error {{'__builtin_amdgcn_global_load_monitor_b32' must be a constant integer}} + *b64out = __builtin_amdgcn_global_load_monitor_b64(b64gaddr, cpol); // expected-error {{'__builtin_amdgcn_global_load_monitor_b64' must be a constant integer}} + *b128out = __builtin_amdgcn_global_load_monitor_b128(b128gaddr, cpol); // expected-error {{'__builtin_amdgcn_global_load_monitor_b128' must be a constant integer}} + *b32out = __builtin_amdgcn_flat_load_monitor_b32(b32faddr, cpol); // expected-error {{'__builtin_amdgcn_flat_load_monitor_b32' must be a constant integer}} + *b64out = __builtin_amdgcn_flat_load_monitor_b64(b64faddr, cpol); // expected-error {{'__builtin_amdgcn_flat_load_monitor_b64' must be a constant integer}} + *b128out = __builtin_amdgcn_flat_load_monitor_b128(b128faddr, cpol); // expected-error {{'__builtin_amdgcn_flat_load_monitor_b128' must be a constant integer}} +} + void test_amdgcn_tensor_load_store(v4i sg0, v8i sg1, v4i sg2, v4i sg3, int cpol) { __builtin_amdgcn_tensor_load_to_lds(sg0, sg1, sg2, sg3, cpol); // expected-error {{'__builtin_amdgcn_tensor_load_to_lds' must be a constant integer}} @@ -36,6 +48,11 @@ void test_amdgcn_tensor_load_store(v4i sg0, v8i sg1, v4i sg2, v4i sg3, int cpol) __builtin_amdgcn_tensor_store_from_lds_d2(sg0, sg1, cpol); // expected-error {{'__builtin_amdgcn_tensor_store_from_lds_d2' must be a constant integer}} } +void test_prefetch(generic void *fptr, global void *gptr, int cpol) { + __builtin_amdgcn_flat_prefetch(fptr, cpol); // expected-error {{'__builtin_amdgcn_flat_prefetch' must be a constant integer}} + __builtin_amdgcn_global_prefetch(gptr, cpol); // expected-error {{'__builtin_amdgcn_global_prefetch' must be a constant integer}} +} + void test_cvt_f32_fp8_e5m3(global int* out, int a) { *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, a); // expected-error {{'__builtin_amdgcn_cvt_f32_fp8_e5m3' must be a constant integer}} diff --git a/clang/test/SemaTemplate/concepts.cpp b/clang/test/SemaTemplate/concepts.cpp index 62a4f95..663bc98 100644 --- a/clang/test/SemaTemplate/concepts.cpp +++ b/clang/test/SemaTemplate/concepts.cpp @@ -1250,3 +1250,11 @@ static_assert(!D<Priv>::has, "Private should be invisible."); static_assert(!D<Prot>::has, "Protected should be invisible."); } + + +namespace GH149986 { +template <typename T> concept PerfectSquare = [](){} // expected-note 2{{here}} +([](auto) { return true; }) < PerfectSquare <class T>; +// expected-error@-1 {{declaration of 'T' shadows template parameter}} \ +// expected-error@-1 {{a concept definition cannot refer to itself}} +} diff --git a/clang/test/SemaTemplate/deduction-guide.cpp b/clang/test/SemaTemplate/deduction-guide.cpp index 0953f647..f6bc6ee 100644 --- a/clang/test/SemaTemplate/deduction-guide.cpp +++ b/clang/test/SemaTemplate/deduction-guide.cpp @@ -966,3 +966,19 @@ Expand<Type, Invocable<>> _{}; // CHECK-NEXT: | `-ParmVarDecl {{.+}} 'T...' pack } + +namespace GH134613 { +template <typename R> struct Foo { + using value_type = R; + + Foo() = default; + Foo(Foo<Foo<R>> &&rhs) {} +}; + +void main() { + auto r1 = Foo(Foo<Foo<int>>{}); + + static_assert(__is_same(decltype(r1)::value_type, int)); +} + +} |