From 3167b3035e8a7fc9c96e00f8e7aad9de218cd0b8 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Fri, 22 Feb 2019 14:42:48 +0000 Subject: [PATCH] [CUDA]Delayed diagnostics for the asm instructions. Adapted targetDiag for the CUDA and used for the delayed diagnostics in asm constructs. Works for both host and device compilation sides. Differential Revision: https://reviews.llvm.org/D58463 llvm-svn: 354671 --- clang/include/clang/Sema/Sema.h | 7 +- clang/lib/Sema/Sema.cpp | 16 +-- clang/lib/Sema/SemaExprCXX.cpp | 2 +- clang/lib/Sema/SemaStmt.cpp | 2 +- clang/lib/Sema/SemaStmtAsm.cpp | 9 -- clang/test/SemaCUDA/asm_delayed_diags.cu | 118 +++++++++++++++++++++++ 6 files changed, 134 insertions(+), 20 deletions(-) create mode 100644 clang/test/SemaCUDA/asm_delayed_diags.cu diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 68730a509a33..94a7aacd382d 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10208,8 +10208,9 @@ public: const T &Value) { if (Diag.ImmediateDiag.hasValue()) *Diag.ImmediateDiag << Value; - else if (Diag.PartialDiag.hasValue()) - *Diag.PartialDiag << Value; + else if (Diag.PartialDiagId.hasValue()) + Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second + << Value; return Diag; } @@ -10223,7 +10224,7 @@ public: // Invariant: At most one of these Optionals has a value. // FIXME: Switch these to a Variant once that exists. llvm::Optional ImmediateDiag; - llvm::Optional PartialDiag; + llvm::Optional PartialDiagId; }; /// Indicate that this function (and thus everything it transtively calls) diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index bb1e264236ff..a0bb0b039c24 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1402,7 +1402,9 @@ Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc, break; case K_Deferred: assert(Fn && "Must have a function to attach the deferred diag to."); - PartialDiag.emplace(S.PDiag(DiagID)); + auto &Diags = S.DeviceDeferredDiags[Fn]; + PartialDiagId.emplace(Diags.size()); + Diags.emplace_back(Loc, S.PDiag(DiagID)); break; } } @@ -1416,9 +1418,9 @@ Sema::DeviceDiagBuilder::~DeviceDiagBuilder() { ImmediateDiag.reset(); // Emit the immediate diag. if (IsWarningOrError && ShowCallStack) emitCallStackNotes(S, Fn); - } else if (PartialDiag) { - assert(ShowCallStack && "Must always show call stack for deferred diags."); - S.DeviceDeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)}); + } else { + assert((!PartialDiagId || ShowCallStack) && + "Must always show call stack for deferred diags."); } } @@ -1487,10 +1489,12 @@ void Sema::markKnownEmitted( } } -Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, - unsigned DiagID) { +Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) { if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) return diagIfOpenMPDeviceCode(Loc, DiagID); + if (getLangOpts().CUDA) + return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID) + : CUDADiagIfHostCode(Loc, DiagID); return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID, getCurFunctionDecl(), *this); } diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index fe1676abb933..512e256d4580 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -750,7 +750,7 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex, bool IsThrownVarInScope) { // Don't report an error if 'throw' is used in system headers. if (!getLangOpts().CXXExceptions && - !getSourceManager().isInSystemHeader(OpLoc)) { + !getSourceManager().isInSystemHeader(OpLoc) && !getLangOpts().CUDA) { // Delay error emission for the OpenMP device code. targetDiag(OpLoc, diag::err_exceptions_disabled) << "throw"; } diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp index fb5c2dacc807..07af218b77ce 100644 --- a/clang/lib/Sema/SemaStmt.cpp +++ b/clang/lib/Sema/SemaStmt.cpp @@ -3993,7 +3993,7 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock, ArrayRef Handlers) { // Don't report an error if 'try' is used in system headers. if (!getLangOpts().CXXExceptions && - !getSourceManager().isInSystemHeader(TryLoc)) { + !getSourceManager().isInSystemHeader(TryLoc) && !getLangOpts().CUDA) { // Delay error emission for the OpenMP device code. targetDiag(TryLoc, diag::err_exceptions_disabled) << "try"; } diff --git a/clang/lib/Sema/SemaStmtAsm.cpp b/clang/lib/Sema/SemaStmtAsm.cpp index 0848ed4ceae9..e1ea3b664631 100644 --- a/clang/lib/Sema/SemaStmtAsm.cpp +++ b/clang/lib/Sema/SemaStmtAsm.cpp @@ -253,15 +253,6 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple, // 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()); diff --git a/clang/test/SemaCUDA/asm_delayed_diags.cu b/clang/test/SemaCUDA/asm_delayed_diags.cu new file mode 100644 index 000000000000..20c5c4480e4d --- /dev/null +++ b/clang/test/SemaCUDA/asm_delayed_diags.cu @@ -0,0 +1,118 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -triple x86_64-unknown-linux-gnu +// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -DHOST_USED -triple x86_64-unknown-linux-gnu +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE_NOT_USED -triple nvptx-unknown-cuda +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -triple nvptx-unknown-cuda +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -DDEVICE_USED -triple nvptx-unknown-cuda + +// REQUIRES: x86-registered-target +// REQUIRES: nvptx-registered-target + +#if (defined(HOST) && !defined(HOST_USED)) || defined(DEVICE_NOT_USED) +// expected-no-diagnostics +#endif + +#include "Inputs/cuda.h" + +static __device__ __host__ void t1(int r) { + __asm__("PR3908 %[lf] %[xx] %[li] %[r]" + : [ r ] "+r"(r) + : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0))); +} + +static __device__ __host__ unsigned t2(signed char input) { + unsigned output; + __asm__("xyz" + : "=a"(output) + : "0"(input)); + return output; +} + +static __device__ __host__ double t3(double x) { + register long double result; + __asm __volatile("frndint" + : "=t"(result) + : "0"(x)); + return result; +} + +static __device__ __host__ unsigned char t4(unsigned char a, unsigned char b) { + unsigned int la = a; + unsigned int lb = b; + unsigned int bigres; + unsigned char res; + __asm__("0:\n1:\n" + : [ bigres ] "=la"(bigres) + : [ la ] "0"(la), [ lb ] "c"(lb) + : "edx", "cc"); + res = bigres; + return res; +} + +static __device__ __host__ void t5(void) { + __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))); +static __device__ __host__ void t6(__m256i *p) { + __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p) + : "ymm0"); +} + +static __device__ __host__ void t7(__m256i *p) { + __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p) + : "r0"); +} + +#ifdef DEVICE +__device__ int m() { + t1(0); + t2(0); + t3(0); + t4(0, 0); + t5(); + t6(0); +#ifdef DEVICE_USED + t7(0); +#endif // DEVICE_USED + return 0; +} +#endif // DEVICE + +#ifdef HOST +__host__ int main() { + t1(0); + t2(0); + t3(0); + t4(0, 0); + t5(); + t6(0); +#ifdef HOST_USED + t7(0); +#endif // HOST_USED + return 0; +} +#endif // HOST + +#if defined(HOST_USED) +// expected-error@69 {{unknown register name 'r0' in asm}} +// expected-note@96 {{called by 'main'}} +#elif defined(DEVICE) +// expected-error@19 {{invalid input constraint 'mx' in asm}} +// expected-error@25 {{invalid output constraint '=a' in asm}} +// expected-error@33 {{invalid output constraint '=t' in asm}} +// expected-error@44 {{invalid output constraint '=la' in asm}} +// expected-error@56 {{unknown register name 'st' in asm}} +// expected-error@64 {{unknown register name 'ymm0' in asm}} +// expected-note@74 {{called by 'm'}} +// expected-note@75 {{called by 'm'}} +// expected-note@76 {{called by 'm'}} +// expected-note@77 {{called by 'm'}} +// expected-note@78 {{called by 'm'}} +// expected-note@79 {{called by 'm'}} +#endif