r354486 - [OPENMP] Delay emission of the asm target-specific error messages.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 20 09:42:57 PST 2019
Author: abataev
Date: Wed Feb 20 09:42:57 2019
New Revision: 354486
URL: http://llvm.org/viewvc/llvm-project?rev=354486&view=rev
Log:
[OPENMP] Delay emission of the asm target-specific error messages.
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.
Reviewers: rjmccall
Subscribers: guansong, jdoerfert, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D58243
Added:
cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/Sema.cpp
cfe/trunk/lib/Sema/SemaStmtAsm.cpp
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=354486&r1=354485&r2=354486&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Feb 20 09:42:57 2019
@@ -10275,6 +10275,8 @@ public:
/// // Otherwise, continue parsing as normal.
DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
+ DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+
enum CUDAFunctionTarget {
CFT_Device,
CFT_Global,
Modified: cfe/trunk/lib/Sema/Sema.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=354486&r1=354485&r2=354486&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/Sema.cpp (original)
+++ cfe/trunk/lib/Sema/Sema.cpp Wed Feb 20 09:42:57 2019
@@ -1487,6 +1487,14 @@ void Sema::markKnownEmitted(
}
}
+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
Modified: cfe/trunk/lib/Sema/SemaStmtAsm.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaStmtAsm.cpp?rev=354486&r1=354485&r2=354486&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaStmtAsm.cpp (original)
+++ cfe/trunk/lib/Sema/SemaStmtAsm.cpp Wed Feb 20 09:42:57 2019
@@ -272,9 +272,9 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
TargetInfo::ConstraintInfo Info(Literal->getString(), OutputName);
if (!Context.getTargetInfo().validateOutputConstraint(Info))
- return StmtError(
- Diag(Literal->getBeginLoc(), diag::err_asm_invalid_output_constraint)
- << Info.getConstraintStr());
+ return StmtResult(targetDiag(Literal->getBeginLoc(),
+ diag::err_asm_invalid_output_constraint)
+ << Info.getConstraintStr());
ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
if (ER.isInvalid())
@@ -327,11 +327,10 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
}
unsigned Size = Context.getTypeSize(OutputExpr->getType());
- if (!Context.getTargetInfo().validateOutputSize(Literal->getString(),
- Size))
- return StmtError(
- Diag(OutputExpr->getBeginLoc(), diag::err_asm_invalid_output_size)
- << Info.getConstraintStr());
+ 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;
@@ -347,9 +346,9 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
TargetInfo::ConstraintInfo Info(Literal->getString(), InputName);
if (!Context.getTargetInfo().validateInputConstraint(OutputConstraintInfos,
Info)) {
- return StmtError(
- Diag(Literal->getBeginLoc(), diag::err_asm_invalid_input_constraint)
- << Info.getConstraintStr());
+ return StmtResult(targetDiag(Literal->getBeginLoc(),
+ diag::err_asm_invalid_input_constraint)
+ << Info.getConstraintStr());
}
ExprResult ER = CheckPlaceholderExpr(Exprs[i]);
@@ -421,8 +420,8 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
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());
}
@@ -434,9 +433,9 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
StringRef Clobber = Literal->getString();
if (!Context.getTargetInfo().isValidClobber(Clobber))
- return StmtError(
- Diag(Literal->getBeginLoc(), diag::err_asm_unknown_register_name)
- << Clobber);
+ return StmtResult(targetDiag(Literal->getBeginLoc(),
+ diag::err_asm_unknown_register_name)
+ << Clobber);
}
GCCAsmStmt *NS =
@@ -447,11 +446,10 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
// 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) {
@@ -488,16 +486,15 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
if (!Context.getTargetInfo().validateConstraintModifier(
Literal->getString(), Piece.getModifier(), Size,
SuggestedModifier)) {
- Diag(Exprs[ConstraintIdx]->getBeginLoc(),
- diag::warn_asm_mismatched_size_modifier);
+ targetDiag(Exprs[ConstraintIdx]->getBeginLoc(),
+ diag::warn_asm_mismatched_size_modifier);
if (!SuggestedModifier.empty()) {
- auto B = Diag(Piece.getRange().getBegin(),
- diag::note_asm_missing_constraint_modifier)
+ 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);
}
}
}
@@ -511,9 +508,10 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
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(),
+ diag::err_asm_unexpected_constraint_alternatives)
+ << NumAlternatives << AltCount);
}
SmallVector<size_t, 4> InputMatchedToOutput(OutputConstraintInfos.size(),
~0U);
@@ -524,9 +522,10 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
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(),
+ diag::err_asm_unexpected_constraint_alternatives)
+ << NumAlternatives << AltCount);
// 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
@@ -541,13 +540,14 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
// 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(),
- diag::err_asm_input_duplicate_match)
- << TiedTo;
- Diag(NS->getInputExpr(InputMatchedToOutput[TiedTo])->getBeginLoc(),
- diag::note_asm_input_duplicate_first)
+ targetDiag(NS->getInputExpr(i)->getBeginLoc(),
+ diag::err_asm_input_duplicate_match)
<< TiedTo;
- return StmtError();
+ return StmtResult(
+ targetDiag(
+ NS->getInputExpr(InputMatchedToOutput[TiedTo])->getBeginLoc(),
+ diag::note_asm_input_duplicate_first)
+ << TiedTo);
}
InputMatchedToOutput[TiedTo] = i;
@@ -632,10 +632,10 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
continue;
}
- Diag(InputExpr->getBeginLoc(), diag::err_asm_tying_incompatible_types)
- << InTy << OutTy << OutputExpr->getSourceRange()
- << InputExpr->getSourceRange();
- return StmtError();
+ return StmtResult(targetDiag(InputExpr->getBeginLoc(),
+ diag::err_asm_tying_incompatible_types)
+ << InTy << OutTy << OutputExpr->getSourceRange()
+ << InputExpr->getSourceRange());
}
// Check for conflicts between clobber list and input or output lists
@@ -643,7 +643,8 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceL
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;
}
Added: cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c?rev=354486&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c (added)
+++ cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c Wed Feb 20 09:42:57 2019
@@ -0,0 +1,118 @@
+// 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 at +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 at +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 at +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 at +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 at +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 at +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 at +2 {{called by 'main'}}
+#endif // DELAYED
+ t1(0);
+#ifdef DELAYED
+// expected-note at +2 {{called by 'main'}}
+#endif // DELAYED
+ t2(0);
+#ifdef DELAYED
+// expected-note at +2 {{called by 'main'}}
+#endif // DELAYED
+ t3(0);
+#ifdef DELAYED
+// expected-note at +2 {{called by 'main'}}
+#endif // DELAYED
+ t4(0, 0);
+#ifdef DELAYED
+// expected-note at +2 {{called by 'main'}}
+#endif // DELAYED
+ t5();
+#ifdef DELAYED
+// expected-note at +2 {{called by 'main'}}
+#endif // DELAYED
+ t6(0);
+ }
+ return 0;
+}
More information about the cfe-commits
mailing list