[OpenMP] Attribute target diagnostics properly

[jdoerfert]

Summary

Type errors in function declarations were not (always) diagnosed prior
to this patch. Furthermore, certain remarks did not get associated
properly which caused them to be emitted multiple times.

[jdoerfert]

Rebase on top of D95903

[jdoerfert]

friendly ping, also D95928.

These were supposed to be backported into 12 :(

[Fznamznon]

clang/lib/Sema/Sema.cpp

1798

So, we assume that lexical context is a function and if it is not, we assume that the value declaration being checked is a function. Can we actually break this assumption?
For example, will something like this work correctly:

struct S {
    __float128 A = 1;
    int bar = 2 * A;
};
#pragma omp declare target
  S s;
#pragma omp end declare target

1835

I was a bit confused by FPTy variable name and already started writing a comment that we already checked a case FunctionProtoType case, then I realized that it is a check for Function*No*Prototype :) . Can we change FPTy to something closer to FunctionNoProtoType ?

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

42–43

So, boo is diagnosed two times, at the point where it actually used and where it is defined, it looks a bit like a duplication isn't it? I think it makes sense to diagnose only usage point with note pointing to the definition.

[jdoerfert]

Thanks for taking a look!

clang/lib/Sema/Sema.cpp

1798

This is diagnosed neither with nor without the patch.
I'm not sure we should either because 2 * A is a constant expression.
We can make the example more complex using a constructor but I'd say that is not related to this patch either way.

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

42–43

It is diagnosed twice, right. I don't think this is a problem of this patch though, see for example line 54 on the left where we emit the boo error 4 times before. Also, take a look at 181 and following in this file. With this patch we reduce the number of times the error is emitted from 3 to 1. It's a bit of a mixed bag but I think this makes it better.

[JonChesterfield]

I agree this is an improvement on the current status and should land.

Reading through it has made me nervous about our current semantic checking. In particular, I'm not sure suppressing error messages from code that ultimately proves to be discarded as the same as not parsing code that ultimately was not required, but that design choice seems to have been made a long time ago. Presumably lazy diagnostics had better engineering properties than lazy parsing and codegen - I'd be interested to see docs/emails about the early design choices for single source target offloading if someone has them on hand.

[jdoerfert]

Rename variable as requested

clang/include/clang/Sema/Sema.h

public:
   ///   realize that the function will be codegen'ed.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in NVPTX device code.
   ///  if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc,
-                                               unsigned DiagID);
+  SemaDiagnosticBuilder
+  diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID, FunctionDecl *FD);
 
   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as host code".
   ///
   /// - If CurContext is a `declare target` function or it is known that the
   /// function is emitted for the host, emits the diagnostics immediately.
   /// - If CurContext is a non-host function, just ignore it.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in NVPTX device code.
   ///  if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
   SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc,
-                                             unsigned DiagID);
+                                             unsigned DiagID, FunctionDecl *FD);
 
-  SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID,
+                                   FunctionDecl *FD = nullptr);
   SemaDiagnosticBuilder targetDiag(SourceLocation Loc,
-                                   const PartialDiagnostic &PD) {
-    return targetDiag(Loc, PD.getDiagID()) << PD;
+                                   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(const ValueDecl *D, SourceLocation Loc);
+  void checkDeviceDecl(ValueDecl *D, SourceLocation Loc);
 
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };

clang/lib/Sema/Sema.cpp

 //===--- Sema.cpp - AST Builder and Semantic Analysis Implementation ------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the actions class which performs semantic analysis and
 // builds an AST out of a parse stream.
 //
 //===----------------------------------------------------------------------===//
 
 #include "UsedDeclVisitor.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/ASTDiagnostic.h"
+#include "clang/AST/Decl.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/PrettyDeclStackTrace.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/Basic/DiagnosticOptions.h"
if (ImmediateDiag) {
     if (IsWarningOrError && ShowCallStack)
       emitCallStackNotes(S, Fn);
   } else {
     assert((!PartialDiagId || ShowCallStack) &&
            "Must always show call stack for deferred diags.");
   }
 }
 
-Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc,
-                                             unsigned DiagID) {
+Sema::SemaDiagnosticBuilder
+Sema::targetDiag(SourceLocation Loc, unsigned DiagID, FunctionDecl *FD) {
+  FD = FD ? FD : getCurFunctionDecl();
   if (LangOpts.OpenMP)
-    return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
-                                   : diagIfOpenMPHostCode(Loc, DiagID);
+    return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID, FD)
+                                   : diagIfOpenMPHostCode(Loc, DiagID, FD);
   if (getLangOpts().CUDA)
     return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
                                       : CUDADiagIfHostCode(Loc, DiagID);
 
   if (getLangOpts().SYCLIsDevice)
     return SYCLDiagIfDeviceCode(Loc, DiagID);
 
   return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc, DiagID,
-                               getCurFunctionDecl(), *this);
+                               FD, *this);
 }
 
 Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID,
                                        bool DeferHint) {
   bool IsError = Diags.getDiagnosticIDs()->isDefaultMappingAsError(DiagID);
   bool ShouldDefer = getLangOpts().CUDA && LangOpts.GPUDeferDiag &&
                      DiagnosticIDs::isDeferrable(DiagID) &&
                      (DeferHint || !IsError);
   auto SetIsLastErrorImmediate = [&](bool Flag) {
     if (IsError)
       IsLastErrorImmediate = Flag;
   };
   if (!ShouldDefer) {
     SetIsLastErrorImmediate(true);
     return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc,
                                  DiagID, getCurFunctionDecl(), *this);
   }
 
-  SemaDiagnosticBuilder DB =
-      getLangOpts().CUDAIsDevice
+  SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice
                                  ? CUDADiagIfDeviceCode(Loc, DiagID)
                                  : CUDADiagIfHostCode(Loc, DiagID);
   SetIsLastErrorImmediate(DB.isImmediate());
   return DB;
 }
 
-void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) {
+void Sema::checkDeviceDecl(ValueDecl *D, SourceLocation Loc) {
   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);
   auto CheckType = [&](QualType Ty) {
     if (Ty->isDependentType())
       return;
 
     if (Ty->isExtIntType()) {
       if (!Context.getTargetInfo().hasExtIntType()) {
-        targetDiag(Loc, diag::err_device_unsupported_type)
+        targetDiag(Loc, diag::err_device_unsupported_type, FD)
             << D << false /*show bit size*/ << 0 /*bitsize*/
             << Ty << Context.getTargetInfo().getTriple().str();
       }
       return;
     }
 
     if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
         ((Ty->isFloat128Type() ||
           (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
          !Context.getTargetInfo().hasFloat128Type()) ||
         (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
          !Context.getTargetInfo().hasInt128Type())) {
-      targetDiag(Loc, diag::err_device_unsupported_type)
+      if (targetDiag(Loc, diag::err_device_unsupported_type, FD)
           << D << true /*show bit size*/
           << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
-          << Context.getTargetInfo().getTriple().str();
-      targetDiag(D->getLocation(), diag::note_defined_here) << 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<FunctionProtoType>(Ty)) {
     for (const auto &ParamTy : FPTy->param_types())
       CheckType(ParamTy);
     CheckType(FPTy->getReturnType());
   }
+  if (const auto *FPTy = dyn_cast<FunctionNoProtoType>(Ty))
+    CheckType(FPTy->getReturnType());
 }
 
 /// Looks through the macro-expansion chain for the given
 /// location, looking for a macro expansion with the given name.
 /// If one is found, returns true and sets the location to that
 /// expansion loc.
 bool Sema::findMacroSpelling(SourceLocation &locref, StringRef name) {
   SourceLocation loc = locref;

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());
+
   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

if (FD && Target != CFT_Host) {
           !VD->getType()->isCUDADeviceBuiltinTextureType() &&
           !VD->isConstexpr() && !VD->getType().isConstQualified())
         targetDiag(*Locs.begin(), diag::err_ref_bad_target)
             << /*host*/ 2 << /*variable*/ 1 << VD << Target;
     }
   }
 
   if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
-    if (const auto *VD = dyn_cast<ValueDecl>(D))
+    if (auto *VD = dyn_cast<ValueDecl>(D))
       checkDeviceDecl(VD, Loc);
 
     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);
   }
 

clang/lib/Sema/SemaOpenMP.cpp

 enum class FunctionEmissionStatus {
   Emitted,
   Discarded,
   Unknown,
 };
 } // anonymous namespace
 
 Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
-                                                         unsigned DiagID) {
+                                                         unsigned DiagID,
+                                                         FunctionDecl *FD) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
 
-  FunctionDecl *FD = getCurFunctionDecl();
   SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
   if (FD) {
     FunctionEmissionStatus FES = getEmissionStatus(FD);
     switch (FES) {
     case FunctionEmissionStatus::Emitted:
       Kind = SemaDiagnosticBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::Unknown:
       Kind = isOpenMPDeviceDelayedContext(*this)
                  ? SemaDiagnosticBuilder::K_Deferred
                  : SemaDiagnosticBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::TemplateDiscarded:
     case FunctionEmissionStatus::OMPDiscarded:
       Kind = SemaDiagnosticBuilder::K_Nop;
       break;
     case FunctionEmissionStatus::CUDADiscarded:
       llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
       break;
     }
   }
 
-  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return SemaDiagnosticBuilder(Kind, Loc, DiagID, FD, *this);
 }
 
 Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
-                                                       unsigned DiagID) {
+                                                       unsigned DiagID,
+                                                       FunctionDecl *FD) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
-  FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
+  FunctionEmissionStatus FES = getEmissionStatus(FD);
   SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
   switch (FES) {
   case FunctionEmissionStatus::Emitted:
     Kind = SemaDiagnosticBuilder::K_Immediate;
     break;
   case FunctionEmissionStatus::Unknown:
     Kind = SemaDiagnosticBuilder::K_Deferred;
     break;
   case FunctionEmissionStatus::TemplateDiscarded:
   case FunctionEmissionStatus::OMPDiscarded:
   case FunctionEmissionStatus::CUDADiscarded:
     Kind = SemaDiagnosticBuilder::K_Nop;
     break;
   }
 
-  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return SemaDiagnosticBuilder(Kind, Loc, DiagID, FD, *this);
 }
 
 static OpenMPDefaultmapClauseKind
 getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
   if (LO.OpenMP <= 45) {
     if (VD->getType().getNonReferenceType()->isScalarType())
       return OMPC_DEFAULTMAP_scalar;
     return OMPC_DEFAULTMAP_aggregate;

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

struct T1 {
   T1() : a(12), f(15) {}
   T1 &operator/(T1 &b) {
     f /= b.a;
     return *this;
   }
 };
 
 #ifndef _ARCH_PPC
-// expected-note@+1 {{'boo' defined here}}
+// 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-note@+1 {{'boo' defined here}}
+// 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
 #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@+4 {{'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 device 'nvptx64-unknown-unknown' does not support it}}
 #else
-// expected-error@+2 {{'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 device 'nvptx64-unknown-unknown' does not support it}}
 #endif
+// expected-note@+1 {{called by 'foo'}}
   boo(0);
   return;
 }
 T bar() {
   return T();
 }
 
 void baz() {
// expected-error@+1 {{'b' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
   a = &b;
 }
 template<bool>
 void dead_template_declare_target() {
   long double *a, b = 0;
   a = &b;
 }
 
-// TODO: We should diagnose the return type and argument type here.
+// 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}}
 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}}
 void ld_arg1a(long double ld) {}
 
 // TODO: We should diagnose the return type and argument type here.
 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}}
 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}}
 void ld_arg1b(ld_ty ld) {}
 
+// TODO: These errors should not be emitted.
+// expected-note@+2 {{'ld_return1c' defined here}}
+// expected-error@+1 {{'ld_return1c' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 static long double ld_return1c() { return 0; }
+// expected-note@+2 {{'ld_arg1c' defined here}}
+// expected-error@+1 {{'ld_arg1c' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 static void ld_arg1c(long double ld) {}
 
+// TODO: These errors should not be emitted.
+// expected-note@+2 {{'ld_return1d' defined here}}
+// expected-error@+1 {{'ld_return1d' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 inline long double ld_return1d() { return 0; }
+// expected-note@+2 {{'ld_arg1d' defined here}}
+// expected-error@+1 {{'ld_arg1d' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 inline void ld_arg1d(long double ld) {}
 
 #ifndef _ARCH_PPC
-// TODO: We should diagnose the return type and argument type here.
+// 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}}
 __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}}
 void ld_arg2a(__float128 ld) {}
 
-// TODO: We should diagnose the return type and argument type here.
 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}}
 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}}
 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 3{{'f' defined here}}
+// expected-note@+1 {{'f' defined here}}
 inline long double dead_inline(long double f) {
 #pragma omp target map(f)
-// TODO: We should not emit the same error message 3 times, here and elsewhere in this file.
-  // expected-error@+1 3{{'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 device '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 3{{'f' defined here}}
+// expected-note@+1 {{'f' defined here}}
 static long double dead_static(long double f) {
 #pragma omp target map(f)
-  // expected-error@+1 3{{'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 device '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 3{{'f' defined here}}
+// expected-note@+1 {{'f' defined here}}
 __float128 foo2(__float128 f) {
 #pragma omp target map(f)
-  // expected-error@+1 3{{'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 device 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 #else
-// expected-note@+1 3{{'f' defined here}}
+// expected-note@+1 {{'f' defined here}}
 long double foo3(long double f) {
 #pragma omp target map(f)
-  // expected-error@+1 3{{'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 device 'nvptx64-unknown-unknown' does not support it}}
   f = 1;
   return f;
 }
 #endif
 
 T foo3() {
   T S;
 #pragma omp target map(S)