[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 ↗

(On Diff #186860)

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

[ABataev]

lib/Sema/SemaStmtAsm.cpp

256–263 ↗

(On Diff #186860)

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 ↗

(On Diff #186860)

SGTM.

[kkwli0]

The change looks okay to me.

cfe/trunk/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
   };
 

cfe/trunk/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;
 

cfe/trunk/lib/Sema/SemaStmtAsm.cpp

for (unsigned i = 0; i != NumOutputs; 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;

cfe/trunk/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;
+}