[clang] Check unsupported types in expressions

[asavonic]

Summary

The patch adds diagnostics for cases like:

float F3 = ((__float128)F1 * (__float128)F2) / 2.0f;

Sema::checkDeviceDecl (renamed to checkTypeSupport) is changed to work
with a type without the corresponding ValueDecl. It is also refactored
so that host diagnostics for unsupported types can be added here as
well.

[Fznamznon]

No objections, just a couple of minors.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10693

Maybe it is just me, but the old version was a bit more clear since it pointed to the wrong declaration. Now looking at the tests I see that clang ends up emitting several identical diagnostics at the same line, but they point to the different declarations/expressions used on this line. Should we still mention which concrete declaration/expression the diagnostic is about?

clang/include/clang/Sema/Sema.h

12163–12165

Having that in single source offloading programming models usually two targets are used, "current target" sounds a bit more clear.

[asavonic]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10693

Agree, but I couldn't figure out good wording for the diagnostic that works well for both decl and expr cases.

[asavonic]

• Reworded the diagnostic.

[asavonic]

Rebased

[jdoerfert]

LG, let's wait for @Fznamznon though.

[Fznamznon]

LGTM, thanks

[asavonic]

This is weird... Buildbot reported that lock/omp_init_lock.c test
crashed with SIGSEGV, but I cannot reproduce this on my machine.

I added a check for a null QualType; not sure if it fixes the issue
with the test.

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_stmt_depends_on_loop_counter : Error<
   "the loop %select{initializer|condition}0 expression depends on the current loop control variable">;
 def err_omp_invariant_dependency : Error<
   "expected loop invariant expression">;
 def err_omp_invariant_or_linear_dependency : Error<
   "expected loop invariant expression or '<invariant1> * %0 + <invariant2>' 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_device_unsupported_type
-    : Error<"%0 requires %select{|%2 bit size}1 %3 type support, but device "
+def err_target_unsupported_type
+    : Error<"%0 requires %select{|%2 bit size}1 %3 type support, but target "
             "'%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">;
 def err_omp_device_type_mismatch : Error<
   "'device_type(%0)' does not match previously specified 'device_type(%1)' for the same declaration">;
 def err_omp_wrong_device_function_call : Error<
   "function with 'device_type(%0)' is not available on %select{device|host}1">;
 def note_omp_marked_device_type_here : Note<"marked as 'device_type(%0)' here">;
 def warn_omp_declare_target_after_first_use : Warning<

clang/include/clang/Sema/Sema.h

public:
   SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID,
                                    FunctionDecl *FD = nullptr);
   SemaDiagnosticBuilder targetDiag(SourceLocation Loc,
                                    const PartialDiagnostic &PD,
                                    FunctionDecl *FD = nullptr) {
     return targetDiag(Loc, PD.getDiagID(), FD) << PD;
   }
 
-  /// Check if the expression is allowed to be used in expressions for the
-  /// offloading devices.
-  void checkDeviceDecl(ValueDecl *D, SourceLocation Loc);
+  /// Check if the type is allowed to be used for the current target.
+  void checkTypeSupport(QualType Ty, SourceLocation Loc,
+                        ValueDecl *D = nullptr);
 
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };

clang/lib/Sema/Sema.cpp

Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID,
 
   SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice
                                  ? CUDADiagIfDeviceCode(Loc, DiagID)
                                  : CUDADiagIfHostCode(Loc, DiagID);
   SetIsLastErrorImmediate(DB.isImmediate());
   return DB;
 }
 
-void Sema::checkDeviceDecl(ValueDecl *D, SourceLocation Loc) {
+void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
+  if (!LangOpts.SYCLIsDevice && !(LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+    return;
+
   if (isUnevaluatedContext())
     return;
 
   Decl *C = cast<Decl>(getCurLexicalContext());
 
   // Memcpy operations for structs containing a member with unsupported type
   // are ok, though.
   if (const auto *MD = dyn_cast<CXXMethodDecl>(C)) {
     if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) &&
         MD->isTrivial())
       return;
 
     if (const auto *Ctor = dyn_cast<CXXConstructorDecl>(MD))
       if (Ctor->isCopyOrMoveConstructor() && Ctor->isTrivial())
         return;
   }
 
   // Try to associate errors with the lexical context, if that is a function, or
   // the value declaration otherwise.
-  FunctionDecl *FD =
-      isa<FunctionDecl>(C) ? cast<FunctionDecl>(C) : dyn_cast<FunctionDecl>(D);
+  FunctionDecl *FD = isa<FunctionDecl>(C) ? cast<FunctionDecl>(C)
+                                          : dyn_cast_or_null<FunctionDecl>(D);
+
   auto CheckType = [&](QualType Ty) {
     if (Ty->isDependentType())
       return;
 
     if (Ty->isExtIntType()) {
       if (!Context.getTargetInfo().hasExtIntType()) {
-        targetDiag(Loc, diag::err_device_unsupported_type, FD)
-            << D << false /*show bit size*/ << 0 /*bitsize*/
+        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*/
             << Ty << Context.getTargetInfo().getTriple().str();
       }
       return;
     }
 
     // Check if we are dealing with two 'long double' but with different
     // semantics.
     bool LongDoubleMismatched = false;
     if (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128) {
       const llvm::fltSemantics &Sem = Context.getFloatTypeSemantics(Ty);
       if (!Ty->isIbm128Type() && !Ty->isFloat128Type() &&
           &Sem != &Context.getTargetInfo().getLongDoubleFormat())
         LongDoubleMismatched = true;
     }
 
     if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
         (Ty->isFloat128Type() && !Context.getTargetInfo().hasFloat128Type()) ||
         (Ty->isIbm128Type() && !Context.getTargetInfo().hasIbm128Type()) ||
         (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
          !Context.getTargetInfo().hasInt128Type()) ||
         LongDoubleMismatched) {
-      if (targetDiag(Loc, diag::err_device_unsupported_type, FD)
-          << D << true /*show bit size*/
+      PartialDiagnostic PD = PDiag(diag::err_target_unsupported_type);
+      if (D)
+        PD << D;
+      else
+        PD << "expression";
+
+      if (targetDiag(Loc, PD, FD)
+          << true /*show bit size*/
           << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
-          << Context.getTargetInfo().getTriple().str())
+          << Context.getTargetInfo().getTriple().str()) {
+        if (D)
           D->setInvalidDecl();
+      }
+      if (D)
         targetDiag(D->getLocation(), diag::note_defined_here, FD) << D;
     }
   };
 
-  QualType Ty = D->getType();
   CheckType(Ty);
 
   if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
     for (const auto &ParamTy : FPTy->param_types())
       CheckType(ParamTy);
     CheckType(FPTy->getReturnType());
   }
   if (const auto *FNPTy = dyn_cast<FunctionNoProtoType>(Ty))

clang/lib/Sema/SemaDecl.cpp

if (getLangOpts().OpenCL) {
     LangAS AddressSpace = NewFD->getReturnType().getAddressSpace();
     if (AddressSpace != LangAS::Default) {
       Diag(NewFD->getLocation(),
            diag::err_opencl_return_value_with_address_space);
       NewFD->setInvalidDecl();
     }
   }
 
-  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
-    checkDeviceDecl(NewFD, D.getBeginLoc());
+  checkTypeSupport(NewFD->getType(), D.getBeginLoc(), NewFD);
 
   if (!getLangOpts().CPlusPlus) {
     // Perform semantic checking on the function declaration.
     if (!NewFD->isInvalidDecl() && NewFD->isMain())
       CheckMain(NewFD, D.getDeclSpec());
 
     if (!NewFD->isInvalidDecl() && NewFD->isMSVCRTEntryPoint())
       CheckMSVCRTEntryPoint(NewFD);

clang/lib/Sema/SemaExpr.cpp

bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
 
   DiagnoseAvailabilityOfDecl(D, Locs, UnknownObjCClass, ObjCPropertyAccess,
                              AvoidPartialAvailabilityChecks, ClassReceiver);
 
   DiagnoseUnusedOfDecl(*this, D, Loc);
 
   diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
 
-  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
-    if (auto *VD = dyn_cast<ValueDecl>(D))
-      checkDeviceDecl(VD, Loc);
+    checkTypeSupport(VD->getType(), Loc, VD);
 
+  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
     if (!Context.getTargetInfo().isTLSSupported())
       if (const auto *VD = dyn_cast<VarDecl>(D))
         if (VD->getTLSKind() != VarDecl::TLS_None)
           targetDiag(*Locs.begin(), diag::err_thread_unsupported);
   }
 
   if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
       !isUnevaluatedContext()) {
if (LHSTy->isImageType() || RHSTy->isImageType() ||
         LHSTy->isSamplerT() || RHSTy->isSamplerT() ||
         LHSTy->isPipeType() || RHSTy->isPipeType() ||
         LHSTy->isBlockPointerType() || RHSTy->isBlockPointerType()) {
       ResultTy = InvalidOperands(OpLoc, LHS, RHS);
       return ExprError();
     }
   }
 
+  checkTypeSupport(LHSExpr->getType(), OpLoc, /*ValueDecl*/ nullptr);
+  checkTypeSupport(RHSExpr->getType(), OpLoc, /*ValueDecl*/ nullptr);
+
   switch (Opc) {
   case BO_Assign:
     ResultTy = CheckAssignmentOperands(LHS.get(), RHS, OpLoc, QualType());
     if (getLangOpts().CPlusPlus &&
         LHS.get()->getObjectKind() != OK_ObjCProperty) {
       VK = LHS.get()->getValueKind();
       OK = LHS.get()->getObjectKind();
     }

clang/test/CodeGen/ibm128-unsupported.c

 // RUN: %clang_cc1 -triple powerpc64le -emit-llvm-bc -fopenmp %s \
 // RUN:   -fopenmp-targets=powerpc64le,x86_64 -o %t-ppc-host.bc
 // RUN: %clang_cc1 -verify -triple x86_64 -aux-triple powerpc64le -fopenmp \
 // RUN:   -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc %s \
 // RUN:   -fsyntax-only
 
 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 device 'x86_64' does not support it}}
+    // expected-error@+1 {{'foo' requires 128 bit size '__ibm128' type support, but target 'x86_64' does not support it}}
     foo(arr[i]);
   }
 }

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

#ifndef _ARCH_PPC
   __float128 f;
 #else
   // expected-note@+1 {{'f' defined here}}
   long double f;
 #endif
   char c;
   T() : a(12), f(15) {}
 #ifndef _ARCH_PPC
-// expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
 #else
-// expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
 #endif
   T &operator+(T &b) {
     f += b.a;
     return *this;
   }
 };
 
 struct T1 {
   char a;
   __int128 f;
   __int128 f1;
   char c;
   T1() : a(12), f(15) {}
   T1 &operator/(T1 &b) {
     f /= b.a;
     return *this;
   }
 };
 
 #ifndef _ARCH_PPC
-// expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
 // expected-note@+1 2{{'boo' defined here}}
 void boo(long double A) { return; }
 #endif
 #pragma omp declare target
 T a = T();
 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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+5 {{'boo' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}}
 #else
-// expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
 #endif
 // expected-note@+1 {{called by 'foo'}}
   boo(0);
   return;
 }
 T bar() {
   return T();
 }

 }
 template<bool>
 void dead_template_declare_target() {
   long double *a, b = 0;
   a = &b;
 }
 
 // expected-note@+2 {{'ld_return1a' defined here}}
-// expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
 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 device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
 void ld_arg1b(ld_ty ld) {}
 
 static long double ld_return1c() { return 0; }
 static void ld_arg1c(long double ld) {}
 
 inline long double ld_return1d() { return 0; }
 inline void ld_arg1d(long double ld) {}
 

 static void ld_use2() {
   long double ld = 0;
   ld += 1;
 }
 
 inline void ld_use3() {
 // expected-note@+1 {{'ld' defined here}}
   long double ld = 0;
-// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
   ld += 1;
 }
 static void ld_use4() {
 // expected-note@+1 {{'ld' defined here}}
   long double ld = 0;
-// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// 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}}
   ld += 1;
 }
 
 void external() {
-// expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   void *p1 = reinterpret_cast<void*>(&ld_return1e);
-// expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   void *p2 = reinterpret_cast<void*>(&ld_arg1e);
-// expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   void *p3 = reinterpret_cast<void*>(&ld_return1f);
-// expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   void *p4 = reinterpret_cast<void*>(&ld_arg1f);
 // TODO: The error message "called by" is not great.
 // expected-note@+1 {{called by 'external'}}
   void *p5 = reinterpret_cast<void*>(&ld_use3);
 // expected-note@+1 {{called by 'external'}}
   void *p6 = reinterpret_cast<void*>(&ld_use4);
 }
 
 #ifndef _ARCH_PPC
 // expected-note@+2 {{'ld_return2a' defined here}}
-// expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target '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 device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target 'nvptx64-unknown-unknown' does not support it}}
 void ld_arg2b(fp128_ty ld) {}
 #endif
 
 #pragma omp end declare target
 
 // TODO: There should not be an error here, dead_inline is never emitted.
 // 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 device 'nvptx64-unknown-unknown' does not support it}}
+  // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 
 // TODO: There should not be an error here, dead_static is never emitted.
 // 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 device 'nvptx64-unknown-unknown' does not support it}}
+  // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 
 template<typename T>
 long double dead_template(long double f) {
 #pragma omp target map(f)
   f = 1;
   return f;
 }
 
 #ifndef _ARCH_PPC
 // 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 device 'nvptx64-unknown-unknown' does not support it}}
+  // expected-error@+1 {{'f' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 #else
 // 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 device 'nvptx64-unknown-unknown' does not support it}}
+  // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 #endif
 
 T foo3() {
   T S;
 #pragma omp target map(S)

clang/test/SemaSYCL/float128.cpp

void host_ok(void) {
   Z<__float128> C;
   C.field1 = A;
 }
 
 void usage() {
   // expected-note@+1 3{{'A' defined here}}
   __float128 A;
   Z<__float128> C;
-  // 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}}
+  // 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}}
   C.field1 = A;
-  // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but device 'spir64' does not support it}}
+  // 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}}
   C.bigfield += 1.0;
 
-  // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+  // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but target '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 device 'spir64' does not support it}}
+    // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}}
     auto BB = A;
-    // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    // 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}}
     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'}}
   foo1();
 }
 
 template <typename t>
 void foo2(){};
 
 // expected-note@+3 {{'P' defined here}}
-// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}}
 // expected-note@+1 2{{'foo' defined here}}
 __float128 foo(__float128 P) { return P; }
 
 template <typename Name, typename Func>
 __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
   // expected-note@+1 5{{called by 'kernel}}
   kernelFunc();
 }
 
 int main() {
   // expected-note@+1 {{'CapturedToDevice' defined here}}
   __float128 CapturedToDevice = 1;
   host_ok();
   kernel<class variables>([=]() {
     decltype(CapturedToDevice) D;
-    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but target 'spir64' does not support it}}
     auto C = CapturedToDevice;
     Z<__float128> S;
-    // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    // 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}}
     S.field1 += 1;
-    // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    // 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}}
     S.field = 1;
   });
 
   kernel<class functions>([=]() {
     // expected-note@+1 2{{called by 'operator()'}}
     usage();
     // expected-note@+1 {{'BBBB' defined here}}
     BIGTY BBBB;
     // expected-note@+3 {{called by 'operator()'}}
-    // 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}}
+    // 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}}
     auto A = foo(BBBB);
   });
 
   kernel<class ok>([=]() {
     Z<__float128> S;
     foo2<__float128>();
     auto A = sizeof(CapturedToDevice);
   });
 
   return 0;
 }

clang/test/SemaSYCL/int128.cpp

void host_ok(void) {
   Z<__int128> C;
   C.field1 = A;
 }
 
 void usage() {
   // expected-note@+1 3{{'A' defined here}}
   __int128 A;
   Z<__int128> C;
-  // 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}}
+  // 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}}
   C.field1 = A;
-  // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__int128') type support, but device 'spir64' does not support it}}
+  // 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}}
   C.bigfield += 1.0;
 
-  // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+  // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but target '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 device 'spir64' does not support it}}
+    // expected-error@+1 {{'A' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}}
     auto BB = A;
-    // expected-error@+1 {{'BB' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+    // 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}}
     BB += 1;
   };
 
   // expected-note@+1 {{called by 'usage'}}
   foo1();
 }
 
 template <typename t>
 void foo2(){};
 
 // expected-note@+3 {{'P' defined here}}
-// expected-error@+2 {{'P' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+// expected-error@+2 {{'P' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}}
 // expected-note@+1 2{{'foo' defined here}}
 __int128 foo(__int128 P) { return P; }
 
 void foobar() {
   // expected-note@+1 {{'operator __int128' defined here}}
   struct X { operator  __int128() const; } x;
   bool a = false;
-  // expected-error@+1 {{'operator __int128' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+  // 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}}
   a = x == __int128(0);
 }
 
 template <typename Name, typename Func>
 __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
   // expected-note@+1 6{{called by 'kernel}}
   kernelFunc();
 }
 
 int main() {
   // expected-note@+1 {{'CapturedToDevice' defined here}}
   __int128 CapturedToDevice = 1;
   host_ok();
   kernel<class variables>([=]() {
     decltype(CapturedToDevice) D;
-    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__int128' type support, but target 'spir64' does not support it}}
     auto C = CapturedToDevice;
     Z<__int128> S;
-    // expected-error@+1 {{'field1' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+    // 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}}
     S.field1 += 1;
-    // expected-error@+1 {{'field' requires 128 bit size '__int128' type support, but device 'spir64' does not support it}}
+    // 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}}
     S.field = 1;
   });
 
   kernel<class functions>([=]() {
     // expected-note@+1 2{{called by 'operator()'}}
     usage();
     // expected-note@+1 {{'BBBB' defined here}}
     BIGTY BBBB;
-    // 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-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-note@+1 1{{called by 'operator()'}}
     auto A = foo(BBBB);
     // expected-note@+1 {{called by 'operator()'}}
     foobar();
   });
 
   kernel<class ok>([=]() {
     Z<__int128> S;