[OPENMP] Delay emission of the asm target-specific error messages.

[ABataev]

Summary

Added the ability to emit target-specific builtin assembler error
messages only in case if the function is really is going to be emitted
for the device.

[ABataev]

Ping!

[tra]

lib/Sema/SemaStmtAsm.cpp

256–263

Now that inline asm errors are delayed, do we still need this check?

[ABataev]

lib/Sema/SemaStmtAsm.cpp

256–263

With this patch, it is going to be delayed only for OpenMP. After the commit, we could extend it for CUDA and remove this code. But not at the moment.

[tra]

lib/Sema/SemaStmtAsm.cpp

256–263

SGTM.

[kkwli0]

The change looks okay to me.

include/clang/Sema/Sema.h

public:
   /// 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.
   DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
 
+  DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 

lib/Sema/Sema.cpp

while (!Worklist.empty()) {
     }
 
     // C.Callee is now known-emitted, so we no longer need to maintain its list
     // of callees in DeviceCallGraph.
     S.DeviceCallGraph.erase(CGIt);
   }
 }
 
+Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc,
+                                         unsigned DiagID) {
+  if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)
+    return diagIfOpenMPDeviceCode(Loc, DiagID);
+  return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID,
+                           getCurFunctionDecl(), *this);
+}
+
 /// 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;
   if (!loc.isMacroID()) return false;
 

lib/Sema/SemaStmtAsm.cpp

StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple,
   StringLiteral *AsmString = cast<StringLiteral>(asmString);
   StringLiteral **Clobbers = reinterpret_cast<StringLiteral**>(clobbers.data());
 
   SmallVector<TargetInfo::ConstraintInfo, 4> OutputConstraintInfos;
 
   // The parser verifies that there is a string literal here.
   assert(AsmString->isAscii());
 
   // If we're compiling CUDA file and function attributes indicate that it's not
   // for this compilation side, skip all the checks.
   if (!DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {
     GCCAsmStmt *NS = new (Context) GCCAsmStmt(
         Context, AsmLoc, IsSimple, IsVolatile, NumOutputs, NumInputs, Names,
         Constraints, Exprs.data(), AsmString, NumClobbers, Clobbers, RParenLoc);
     return NS;
   }
 
   for (unsigned i = 0; i != NumOutputs; i++) {
     StringLiteral *Literal = Constraints[i];
     assert(Literal->isAscii());
 
     StringRef OutputName;
     if (Names[i])
       OutputName = Names[i]->getName();
 
     TargetInfo::ConstraintInfo Info(Literal->getString(), OutputName);
     if (!Context.getTargetInfo().validateOutputConstraint(Info))
-      return StmtError(
-          Diag(Literal->getBeginLoc(), diag::err_asm_invalid_output_constraint)
+      return StmtResult(targetDiag(Literal->getBeginLoc(),
+                                   diag::err_asm_invalid_output_constraint)
                         << Info.getConstraintStr());
 
     ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
     if (ER.isInvalid())
       return StmtError();
     Exprs[i] = ER.get();
 
     // Check that the output exprs are valid lvalues.
     Expr *OutputExpr = Exprs[i];
case Expr::MLV_IncompleteVoidType:
       LLVM_FALLTHROUGH;
     default:
       return StmtError(Diag(OutputExpr->getBeginLoc(),
                             diag::err_asm_invalid_lvalue_in_output)
                        << OutputExpr->getSourceRange());
     }
 
     unsigned Size = Context.getTypeSize(OutputExpr->getType());
-    if (!Context.getTargetInfo().validateOutputSize(Literal->getString(),
-                                                    Size))
-      return StmtError(
-          Diag(OutputExpr->getBeginLoc(), diag::err_asm_invalid_output_size)
+    if (!Context.getTargetInfo().validateOutputSize(Literal->getString(), Size))
+      return StmtResult(targetDiag(OutputExpr->getBeginLoc(),
+                                   diag::err_asm_invalid_output_size)
                         << Info.getConstraintStr());
   }
 
   SmallVector<TargetInfo::ConstraintInfo, 4> InputConstraintInfos;
 
   for (unsigned i = NumOutputs, e = NumOutputs + NumInputs; i != e; i++) {
     StringLiteral *Literal = Constraints[i];
     assert(Literal->isAscii());
 
     StringRef InputName;
     if (Names[i])
       InputName = Names[i]->getName();
 
     TargetInfo::ConstraintInfo Info(Literal->getString(), InputName);
     if (!Context.getTargetInfo().validateInputConstraint(OutputConstraintInfos,
                                                          Info)) {
-      return StmtError(
-          Diag(Literal->getBeginLoc(), diag::err_asm_invalid_input_constraint)
+      return StmtResult(targetDiag(Literal->getBeginLoc(),
+                                   diag::err_asm_invalid_input_constraint)
                         << Info.getConstraintStr());
     }
 
     ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
     if (ER.isInvalid())
       return StmtError();
     Exprs[i] = ER.get();
 
     Expr *InputExpr = Exprs[i];
for (unsigned i = NumOutputs, e = NumOutputs + NumInputs; i != e; i++) {
     if (!Ty->isVoidType() || !Info.allowsMemory())
       if (RequireCompleteType(InputExpr->getBeginLoc(), Exprs[i]->getType(),
                               diag::err_dereference_incomplete_type))
         return StmtError();
 
     unsigned Size = Context.getTypeSize(Ty);
     if (!Context.getTargetInfo().validateInputSize(Literal->getString(),
                                                    Size))
-      return StmtError(
-          Diag(InputExpr->getBeginLoc(), diag::err_asm_invalid_input_size)
+      return StmtResult(
+          targetDiag(InputExpr->getBeginLoc(), diag::err_asm_invalid_input_size)
           << Info.getConstraintStr());
   }
 
   // Check that the clobbers are valid.
   for (unsigned i = 0; i != NumClobbers; i++) {
     StringLiteral *Literal = Clobbers[i];
     assert(Literal->isAscii());
 
     StringRef Clobber = Literal->getString();
 
     if (!Context.getTargetInfo().isValidClobber(Clobber))
-      return StmtError(
-          Diag(Literal->getBeginLoc(), diag::err_asm_unknown_register_name)
+      return StmtResult(targetDiag(Literal->getBeginLoc(),
+                                   diag::err_asm_unknown_register_name)
                         << Clobber);
   }
 
   GCCAsmStmt *NS =
     new (Context) GCCAsmStmt(Context, AsmLoc, IsSimple, IsVolatile, NumOutputs,
                              NumInputs, Names, Constraints, Exprs.data(),
                              AsmString, NumClobbers, Clobbers, RParenLoc);
   // Validate the asm string, ensuring it makes sense given the operands we
   // have.
   SmallVector<GCCAsmStmt::AsmStringPiece, 8> Pieces;
   unsigned DiagOffs;
-  if (unsigned DiagID = NS->AnalyzeAsmString(Pieces, Context, DiagOffs)) {
-    Diag(getLocationOfStringLiteralByte(AsmString, DiagOffs), DiagID)
-           << AsmString->getSourceRange();
-    return StmtError();
-  }
+  if (unsigned DiagID = NS->AnalyzeAsmString(Pieces, Context, DiagOffs))
+    return StmtResult(
+        targetDiag(getLocationOfStringLiteralByte(AsmString, DiagOffs), DiagID)
+        << AsmString->getSourceRange());
 
   // Validate constraints and modifiers.
   for (unsigned i = 0, e = Pieces.size(); i != e; ++i) {
     GCCAsmStmt::AsmStringPiece &Piece = Pieces[i];
     if (!Piece.isOperand()) continue;
 
     // Look for the correct constraint index.
     unsigned ConstraintIdx = Piece.getOperandNo();
for (unsigned i = 0, e = Pieces.size(); i != e; ++i) {
     if (Ty->isDependentType() || Ty->isIncompleteType())
       continue;
 
     unsigned Size = Context.getTypeSize(Ty);
     std::string SuggestedModifier;
     if (!Context.getTargetInfo().validateConstraintModifier(
             Literal->getString(), Piece.getModifier(), Size,
             SuggestedModifier)) {
-      Diag(Exprs[ConstraintIdx]->getBeginLoc(),
+      targetDiag(Exprs[ConstraintIdx]->getBeginLoc(),
                  diag::warn_asm_mismatched_size_modifier);
 
       if (!SuggestedModifier.empty()) {
-        auto B = Diag(Piece.getRange().getBegin(),
+        auto B = targetDiag(Piece.getRange().getBegin(),
                             diag::note_asm_missing_constraint_modifier)
                  << SuggestedModifier;
         SuggestedModifier = "%" + SuggestedModifier + Piece.getString();
-        B.AddFixItHint(FixItHint::CreateReplacement(Piece.getRange(),
-                                                    SuggestedModifier));
+        B << FixItHint::CreateReplacement(Piece.getRange(), SuggestedModifier);
       }
     }
   }
 
   // Validate tied input operands for type mismatches.
   unsigned NumAlternatives = ~0U;
   for (unsigned i = 0, e = OutputConstraintInfos.size(); i != e; ++i) {
     TargetInfo::ConstraintInfo &Info = OutputConstraintInfos[i];
     StringRef ConstraintStr = Info.getConstraintStr();
     unsigned AltCount = ConstraintStr.count(',') + 1;
     if (NumAlternatives == ~0U)
       NumAlternatives = AltCount;
     else if (NumAlternatives != AltCount)
-      return StmtError(Diag(NS->getOutputExpr(i)->getBeginLoc(),
-                            diag::err_asm_unexpected_constraint_alternatives)
-                       << NumAlternatives << AltCount);
+      return StmtResult(
+          targetDiag(NS->getOutputExpr(i)->getBeginLoc(),
   }
   SmallVector<size_t, 4> InputMatchedToOutput(OutputConstraintInfos.size(),
                                               ~0U);
   for (unsigned i = 0, e = InputConstraintInfos.size(); i != e; ++i) {
     TargetInfo::ConstraintInfo &Info = InputConstraintInfos[i];
     StringRef ConstraintStr = Info.getConstraintStr();
     unsigned AltCount = ConstraintStr.count(',') + 1;
     if (NumAlternatives == ~0U)
       NumAlternatives = AltCount;
     else if (NumAlternatives != AltCount)
-      return StmtError(Diag(NS->getInputExpr(i)->getBeginLoc(),
-                            diag::err_asm_unexpected_constraint_alternatives)
-                       << NumAlternatives << AltCount);
+      return StmtResult(
+          targetDiag(NS->getInputExpr(i)->getBeginLoc(),
 
     // If this is a tied constraint, verify that the output and input have
     // either exactly the same type, or that they are int/ptr operands with the
     // same size (int/long, int*/long, are ok etc).
     if (!Info.hasTiedOperand()) continue;
 
     unsigned TiedTo = Info.getTiedOperand();
     unsigned InputOpNo = i+NumOutputs;
     Expr *OutputExpr = Exprs[TiedTo];
     Expr *InputExpr = Exprs[InputOpNo];
 
     // Make sure no more than one input constraint matches each output.
     assert(TiedTo < InputMatchedToOutput.size() && "TiedTo value out of range");
     if (InputMatchedToOutput[TiedTo] != ~0U) {
-      Diag(NS->getInputExpr(i)->getBeginLoc(),
+      targetDiag(NS->getInputExpr(i)->getBeginLoc(),
                  diag::err_asm_input_duplicate_match)
           << TiedTo;
-      Diag(NS->getInputExpr(InputMatchedToOutput[TiedTo])->getBeginLoc(),
+      return StmtResult(
+          targetDiag(
+              NS->getInputExpr(InputMatchedToOutput[TiedTo])->getBeginLoc(),
               diag::note_asm_input_duplicate_first)
-          << TiedTo;
-      return StmtError();
+          << TiedTo);
     }
     InputMatchedToOutput[TiedTo] = i;
 
     if (OutputExpr->isTypeDependent() || InputExpr->isTypeDependent())
       continue;
 
     QualType InTy = InputExpr->getType();
     QualType OutTy = OutputExpr->getType();
if (InputDomain == AD_Int && OutputDomain == AD_Int &&
       CastKind castKind =
         (OutTy->isBooleanType() ? CK_IntegralToBoolean : CK_IntegralCast);
       InputExpr = ImpCastExprToType(InputExpr, OutTy, castKind).get();
       Exprs[InputOpNo] = InputExpr;
       NS->setInputExpr(i, InputExpr);
       continue;
     }
 
-    Diag(InputExpr->getBeginLoc(), diag::err_asm_tying_incompatible_types)
+    return StmtResult(targetDiag(InputExpr->getBeginLoc(),
+                                 diag::err_asm_tying_incompatible_types)
                       << InTy << OutTy << OutputExpr->getSourceRange()
-        << InputExpr->getSourceRange();
-    return StmtError();
+                      << InputExpr->getSourceRange());
   }
 
   // Check for conflicts between clobber list and input or output lists
   SourceLocation ConstraintLoc =
       getClobberConflictLocation(Exprs, Constraints, Clobbers, NumClobbers,
                                  Context.getTargetInfo(), Context);
   if (ConstraintLoc.isValid())
-    return Diag(ConstraintLoc, diag::error_inoutput_conflict_with_clobber);
+    return StmtResult(
+        targetDiag(ConstraintLoc, diag::error_inoutput_conflict_with_clobber));
 
   return NS;
 }
 
 void Sema::FillInlineAsmIdentifierInfo(Expr *Res,
                                        llvm::InlineAsmIdentifierInfo &Info) {
   QualType T = Res->getType();
   Expr::EvalResult Eval;

test/OpenMP/nvptx_asm_delayed_diags.c

+// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+#ifndef DIAGS
+// expected-no-diagnostics
+#endif // DIAGS
+
+#ifdef IMMEDIATE
+#pragma omp declare target
+#endif //IMMEDIATE
+void t1(int r) {
+#ifdef DIAGS
+// expected-error@+4 {{invalid input constraint 'mx' in asm}}
+#endif // DIAGS
+  __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
+          : [ r ] "+r"(r)
+          : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
+}
+
+unsigned t2(signed char input) {
+  unsigned output;
+#ifdef DIAGS
+// expected-error@+3 {{invalid output constraint '=a' in asm}}
+#endif // DIAGS
+  __asm__("xyz"
+          : "=a"(output)
+          : "0"(input));
+  return output;
+}
+
+double t3(double x) {
+  register long double result;
+#ifdef DIAGS
+// expected-error@+3 {{invalid output constraint '=t' in asm}}
+#endif // DIAGS
+  __asm __volatile("frndint"
+                   : "=t"(result)
+                   : "0"(x));
+  return result;
+}
+
+unsigned char t4(unsigned char a, unsigned char b) {
+  unsigned int la = a;
+  unsigned int lb = b;
+  unsigned int bigres;
+  unsigned char res;
+#ifdef DIAGS
+// expected-error@+3 {{invalid output constraint '=la' in asm}}
+#endif // DIAGS
+  __asm__("0:\n1:\n"
+          : [ bigres ] "=la"(bigres)
+          : [ la ] "0"(la), [ lb ] "c"(lb)
+          : "edx", "cc");
+  res = bigres;
+  return res;
+}
+
+void t5(void) {
+#ifdef DIAGS
+// expected-error@+6 {{unknown register name 'st' in asm}}
+#endif // DIAGS
+  __asm__ __volatile__(
+      "finit"
+      :
+      :
+      : "st", "st(1)", "st(2)", "st(3)",
+        "st(4)", "st(5)", "st(6)", "st(7)",
+        "fpsr", "fpcr");
+}
+
+typedef long long __m256i __attribute__((__vector_size__(32)));
+void t6(__m256i *p) {
+#ifdef DIAGS
+// expected-error@+3 {{unknown register name 'ymm0' in asm}}
+#endif // DIAGS
+  __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
+                   : "ymm0");
+}
+#ifdef IMMEDIATE
+#pragma omp end declare target
+#endif //IMMEDIATE
+
+int main() {
+#ifdef DELAYED
+#pragma omp target
+#endif // DELAYED
+  {
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t1(0);
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t2(0);
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t3(0);
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t4(0, 0);
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t5();
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+    t6(0);
+  }
+  return 0;
+}