[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

@@ -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<SemaDiagnosticBuilder> ImmediateDiag;
-    llvm::Optional<PartialDiagnostic> PartialDiag;
+    llvm::Optional<unsigned> PartialDiagId;
   };
 
   /// Indicate that this function (and thus everything it transtively calls)

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);
 }

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";
   }

clang/lib/Sema/SemaStmt.cpp

@@ -3993,7 +3993,7 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
                                   ArrayRef<Stmt *> 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";
   }

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());

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