[CUDA]Delayed diagnostics for the asm instructions.

Summary:
Adapted targetDiag for the CUDA and used for the delayed diagnostics in
asm constructs. Works for both host and device compilation sides.

Reviewers: tra, jlebar

Subscribers: jdoerfert, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D58463

llvm-svn: 354593

5 files changed, 124 insertions, 13 deletions

diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index bb1e264236f..de107594eff 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1487,10 +1487,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 fe1676abb93..512e256d458 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 fb5c2dacc80..07af218b77c 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<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";
   }
diff --git a/clang/lib/Sema/SemaStmtAsm.cpp b/clang/lib/Sema/SemaStmtAsm.cpp
index 0848ed4ceae..e1ea3b66463 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 00000000000..20c5c4480e4
--- /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