Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -10269,6 +10269,8 @@ /// // Otherwise, continue parsing as normal. DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID); + DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID); + enum CUDAFunctionTarget { CFT_Device, CFT_Global, Index: lib/Sema/Sema.cpp =================================================================== --- lib/Sema/Sema.cpp +++ lib/Sema/Sema.cpp @@ -1487,6 +1487,14 @@ } } +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 Index: lib/Sema/SemaStmtAsm.cpp =================================================================== --- lib/Sema/SemaStmtAsm.cpp +++ lib/Sema/SemaStmtAsm.cpp @@ -272,9 +272,9 @@ 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 @@ } 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 InputConstraintInfos; @@ -347,9 +346,9 @@ 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 @@ 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 @@ 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 @@ // have. SmallVector 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 @@ 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 @@ 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 InputMatchedToOutput(OutputConstraintInfos.size(), ~0U); @@ -524,9 +522,10 @@ 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 @@ // 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 @@ 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 @@ 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; } Index: test/OpenMP/nvptx_asm_delayed_diags.c =================================================================== --- /dev/null +++ test/OpenMP/nvptx_asm_delayed_diags.c @@ -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@+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; +}