diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9c739c3c1f2506..824d9bf4693607 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10688,8 +10688,8 @@ def err_omp_invariant_or_linear_dependency : Error< "expected loop invariant expression or ' * %0 + ' kind of expression">; def err_omp_wrong_dependency_iterator_type : Error< "expected an integer or a pointer type of the outer loop counter '%0' for non-rectangular nests">; -def err_target_unsupported_type - : Error<"%0 requires %select{|%2 bit size}1 %3 type support, but target " +def err_device_unsupported_type + : Error<"%0 requires %select{|%2 bit size}1 %3 type support, but device " "'%4' does not support it">; def err_omp_lambda_capture_in_declare_target_not_to : Error< "variable captured in declare target region must appear in a to clause">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 9e4175b6de8115..710abeb1ea5142 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12160,9 +12160,9 @@ class Sema final { return targetDiag(Loc, PD.getDiagID(), FD) << PD; } - /// Check if the type is allowed to be used for the current target. - void checkTypeSupport(QualType Ty, SourceLocation Loc, - ValueDecl *D = nullptr); + /// Check if the expression is allowed to be used in expressions for the + /// offloading devices. + void checkDeviceDecl(ValueDecl *D, SourceLocation Loc); enum CUDAFunctionTarget { CFT_Device, diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index b23d92d1010741..a9867697a4c31f 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1852,10 +1852,7 @@ Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID, return DB; } -void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) { - if (!LangOpts.SYCLIsDevice && !(LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) - return; - +void Sema::checkDeviceDecl(ValueDecl *D, SourceLocation Loc) { if (isUnevaluatedContext()) return; @@ -1875,22 +1872,16 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) { // Try to associate errors with the lexical context, if that is a function, or // the value declaration otherwise. - FunctionDecl *FD = isa(C) ? cast(C) - : dyn_cast_or_null(D); - + FunctionDecl *FD = + isa(C) ? cast(C) : dyn_cast(D); auto CheckType = [&](QualType Ty) { if (Ty->isDependentType()) return; if (Ty->isExtIntType()) { if (!Context.getTargetInfo().hasExtIntType()) { - PartialDiagnostic PD = PDiag(diag::err_target_unsupported_type); - if (D) - PD << D; - else - PD << "expression"; - targetDiag(Loc, PD, FD) - << false /*show bit size*/ << 0 /*bitsize*/ + targetDiag(Loc, diag::err_device_unsupported_type, FD) + << D << false /*show bit size*/ << 0 /*bitsize*/ << Ty << Context.getTargetInfo().getTriple().str(); } return; @@ -1912,24 +1903,16 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) { (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 && !Context.getTargetInfo().hasInt128Type()) || LongDoubleMismatched) { - PartialDiagnostic PD = PDiag(diag::err_target_unsupported_type); - if (D) - PD << D; - else - PD << "expression"; - - if (targetDiag(Loc, PD, FD) - << true /*show bit size*/ + if (targetDiag(Loc, diag::err_device_unsupported_type, FD) + << D << true /*show bit size*/ << static_cast(Context.getTypeSize(Ty)) << Ty - << Context.getTargetInfo().getTriple().str()) { - if (D) - D->setInvalidDecl(); - } - if (D) - targetDiag(D->getLocation(), diag::note_defined_here, FD) << D; + << Context.getTargetInfo().getTriple().str()) + D->setInvalidDecl(); + targetDiag(D->getLocation(), diag::note_defined_here, FD) << D; } }; + QualType Ty = D->getType(); CheckType(Ty); if (const auto *FPTy = dyn_cast(Ty)) { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 34af164d10a47f..cc3417d4ccba20 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -9569,7 +9569,8 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC, } } - checkTypeSupport(NewFD->getType(), D.getBeginLoc(), NewFD); + if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) + checkDeviceDecl(NewFD, D.getBeginLoc()); if (!getLangOpts().CPlusPlus) { // Perform semantic checking on the function declaration. diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index cbfe5ae1a588fc..7c8d753730a0a4 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -366,10 +366,10 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc); - if (auto *VD = dyn_cast(D)) - checkTypeSupport(VD->getType(), Loc, VD); - if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) { + if (auto *VD = dyn_cast(D)) + checkDeviceDecl(VD, Loc); + if (!Context.getTargetInfo().isTLSSupported()) if (const auto *VD = dyn_cast(D)) if (VD->getTLSKind() != VarDecl::TLS_None) @@ -14161,9 +14161,6 @@ ExprResult Sema::CreateBuiltinBinOp(SourceLocation OpLoc, } } - checkTypeSupport(LHSExpr->getType(), OpLoc, /*ValueDecl*/ nullptr); - checkTypeSupport(RHSExpr->getType(), OpLoc, /*ValueDecl*/ nullptr); - switch (Opc) { case BO_Assign: ResultTy = CheckAssignmentOperands(LHS.get(), RHS, OpLoc, QualType()); diff --git a/clang/test/CodeGen/ibm128-unsupported.c b/clang/test/CodeGen/ibm128-unsupported.c index 4c4930f3a7ce44..4ade7ddaa1e39a 100644 --- a/clang/test/CodeGen/ibm128-unsupported.c +++ b/clang/test/CodeGen/ibm128-unsupported.c @@ -9,7 +9,7 @@ void foo(__ibm128 x); // expected-note {{'foo' defined here}} void loop(int n, __ibm128 *arr) { #pragma omp target parallel for (int i = 0; i < n; ++i) { - // expected-error@+1 {{'foo' requires 128 bit size '__ibm128' type support, but target 'x86_64' does not support it}} + // expected-error@+1 {{'foo' requires 128 bit size '__ibm128' type support, but device 'x86_64' does not support it}} foo(arr[i]); } } diff --git a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp index e57ef1d288a320..a319c78f73c567 100644 --- a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp @@ -16,11 +16,9 @@ struct T { char c; T() : a(12), f(15) {} #ifndef _ARCH_PPC -// expected-error@+7 {{'f' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} -// expected-error@+6 {{expression requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} #else -// expected-error@+4 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} -// expected-error@+3 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} #endif T &operator+(T &b) { f += b.a; @@ -41,11 +39,11 @@ struct T1 { }; #ifndef _ARCH_PPC -// expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 2{{'boo' defined here}} void boo(__float128 A) { return; } #else -// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 2{{'boo' defined here}} void boo(long double A) { return; } #endif @@ -55,9 +53,9 @@ T f = a; void foo(T a = T()) { a = a + f; // expected-note {{called by 'foo'}} #ifndef _ARCH_PPC -// expected-error@+5 {{'boo' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+5 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} #else -// expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} #endif // expected-note@+1 {{called by 'foo'}} boo(0); @@ -98,18 +96,18 @@ void dead_template_declare_target() { } // expected-note@+2 {{'ld_return1a' defined here}} -// expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} long double ld_return1a() { return 0; } // expected-note@+2 {{'ld_arg1a' defined here}} -// expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg1a(long double ld) {} typedef long double ld_ty; // expected-note@+2 {{'ld_return1b' defined here}} -// expected-error@+1 {{'ld_return1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}} ld_ty ld_return1b() { return 0; } // expected-note@+2 {{'ld_arg1b' defined here}} -// expected-error@+1 {{'ld_arg1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg1b(ld_ty ld) {} static long double ld_return1c() { return 0; } @@ -140,26 +138,24 @@ static void ld_use2() { inline void ld_use3() { // expected-note@+1 {{'ld' defined here}} long double ld = 0; -// expected-error@+2 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} -// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} ld += 1; } static void ld_use4() { // expected-note@+1 {{'ld' defined here}} long double ld = 0; -// expected-error@+2 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} -// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} ld += 1; } void external() { -// expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p1 = reinterpret_cast(&ld_return1e); -// expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p2 = reinterpret_cast(&ld_arg1e); -// expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p3 = reinterpret_cast(&ld_return1f); -// expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p4 = reinterpret_cast(&ld_arg1f); // TODO: The error message "called by" is not great. // expected-note@+1 {{called by 'external'}} @@ -170,18 +166,18 @@ void external() { #ifndef _ARCH_PPC // expected-note@+2 {{'ld_return2a' defined here}} -// expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} __float128 ld_return2a() { return 0; } // expected-note@+2 {{'ld_arg2a' defined here}} -// expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg2a(__float128 ld) {} typedef __float128 fp128_ty; // expected-note@+2 {{'ld_return2b' defined here}} -// expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}} fp128_ty ld_return2b() { return 0; } // expected-note@+2 {{'ld_arg2b' defined here}} -// expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target 'nvptx64-unknown-unknown' does not support it}} +// expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg2b(fp128_ty ld) {} #endif @@ -191,7 +187,7 @@ void ld_arg2b(fp128_ty ld) {} // expected-note@+1 {{'f' defined here}} inline long double dead_inline(long double f) { #pragma omp target map(f) - // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } @@ -200,7 +196,7 @@ inline long double dead_inline(long double f) { // expected-note@+1 {{'f' defined here}} static long double dead_static(long double f) { #pragma omp target map(f) - // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } @@ -216,7 +212,7 @@ long double dead_template(long double f) { // expected-note@+1 {{'f' defined here}} __float128 foo2(__float128 f) { #pragma omp target map(f) - // expected-error@+1 {{'f' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} + // expected-error@+1 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } @@ -224,7 +220,7 @@ __float128 foo2(__float128 f) { // expected-note@+1 {{'f' defined here}} long double foo3(long double f) { #pragma omp target map(f) - // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index 4629437afa4e29..b91535eda4897c 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -26,28 +26,20 @@ void usage() { // expected-note@+1 3{{'A' defined here}} __float128 A; Z<__float128> C; - // expected-error@+3 2{{expression requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - // expected-error@+2 {{'A' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+2 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} C.field1 = A; - // expected-error@+2 {{expression requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but device 'spir64' does not support it}} C.bigfield += 1.0; - // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} auto foo1 = [=]() { __float128 AA; // expected-note@+2 {{'BB' defined here}} - // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} auto BB = A; - // expected-error@+2 {{expression requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} BB += 1; - - float F1 = 0.1f; - float F2 = 0.1f; - // expected-error@+1 3{{expression requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - float F3 = ((__float128)F1 * (__float128)F2) / 2.0f; }; // expected-note@+1 {{called by 'usage'}} @@ -58,7 +50,7 @@ template void foo2(){}; // expected-note@+3 {{'P' defined here}} -// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} +// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} // expected-note@+1 2{{'foo' defined here}} __float128 foo(__float128 P) { return P; } @@ -74,14 +66,12 @@ int main() { host_ok(); kernel([=]() { decltype(CapturedToDevice) D; - // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} auto C = CapturedToDevice; Z<__float128> S; - // expected-error@+2 {{expression requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} S.field1 += 1; - // expected-error@+2 {{expression requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} S.field = 1; }); @@ -91,8 +81,8 @@ int main() { // expected-note@+1 {{'BBBB' defined here}} BIGTY BBBB; // expected-note@+3 {{called by 'operator()'}} - // expected-error@+2 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but target 'spir64' does not support it}} - // expected-error@+1 2{{'foo' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}} + // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}} + // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}} auto A = foo(BBBB); }); diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp index bfa9a041880275..f6f92c237a9c5e 100644 --- a/clang/test/SemaSYCL/int128.cpp +++ b/clang/test/SemaSYCL/int128.cpp @@ -26,22 +26,19 @@ void usage() { // expected-note@+1 3{{'A' defined here}} __int128 A; Z<__int128> C; - // expected-error@+3 2{{expression requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+2 {{'A' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+2 {{'A' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} + // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} C.field1 = A; - // expected-error@+2 {{expression requires 128 bit size 'Z::BIGTYPE' (aka '__int128') type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__int128') type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__int128') type support, but device 'spir64' does not support it}} C.bigfield += 1.0; - // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} auto foo1 = [=]() { __int128 AA; // expected-note@+2 {{'BB' defined here}} - // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} auto BB = A; - // expected-error@+2 {{expression requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'BB' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'BB' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} BB += 1; }; @@ -53,7 +50,7 @@ template void foo2(){}; // expected-note@+3 {{'P' defined here}} -// expected-error@+2 {{'P' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} +// expected-error@+2 {{'P' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} // expected-note@+1 2{{'foo' defined here}} __int128 foo(__int128 P) { return P; } @@ -61,8 +58,7 @@ void foobar() { // expected-note@+1 {{'operator __int128' defined here}} struct X { operator __int128() const; } x; bool a = false; - // expected-error@+2 2{{expression requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'operator __int128' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'operator __int128' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} a = x == __int128(0); } @@ -78,14 +74,12 @@ int main() { host_ok(); kernel([=]() { decltype(CapturedToDevice) D; - // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} auto C = CapturedToDevice; Z<__int128> S; - // expected-error@+2 {{expression requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} S.field1 += 1; - // expected-error@+2 {{expression requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} - // expected-error@+1 {{'field' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+1 {{'field' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} S.field = 1; }); @@ -94,8 +88,8 @@ int main() { usage(); // expected-note@+1 {{'BBBB' defined here}} BIGTY BBBB; - // expected-error@+3 {{'BBBB' requires 128 bit size 'BIGTY' (aka 'unsigned __int128') type support, but target 'spir64' does not support it}} - // expected-error@+2 2{{'foo' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}} + // expected-error@+3 {{'BBBB' requires 128 bit size 'BIGTY' (aka 'unsigned __int128') type support, but device 'spir64' does not support it}} + // expected-error@+2 2{{'foo' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}} // expected-note@+1 1{{called by 'operator()'}} auto A = foo(BBBB); // expected-note@+1 {{called by 'operator()'}}