From: Alexey Bataev <a.bataev@hotmail.com>
Date: Fri, 22 Feb 2019 20:36:10 +0000 (+0000)
Subject: [OPENMP] Delayed diagnostics for VLA support.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a41469a23c44af02ae43941e9534ebc5b3dde54b;p=clang

[OPENMP] Delayed diagnostics for VLA support.

Generalized processing of the deferred diagnostics for OpenMP/CUDA code.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@354690 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index 94a7aacd38..eb5126dc3d 100644
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -10189,6 +10189,8 @@ public:
 
     DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
                       FunctionDecl *Fn, Sema &S);
+    DeviceDiagBuilder(DeviceDiagBuilder &&D);
+    DeviceDiagBuilder(const DeviceDiagBuilder &) = default;
     ~DeviceDiagBuilder();
 
     /// Convertible to bool: True if we immediately emitted an error, false if
diff --git a/lib/Sema/Sema.cpp b/lib/Sema/Sema.cpp
index a0bb0b039c..da87e15683 100644
--- a/lib/Sema/Sema.cpp
+++ b/lib/Sema/Sema.cpp
@@ -1409,6 +1409,16 @@ Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc,
   }
 }
 
+Sema::DeviceDiagBuilder::DeviceDiagBuilder(DeviceDiagBuilder &&D)
+    : S(D.S), Loc(D.Loc), DiagID(D.DiagID), Fn(D.Fn),
+      ShowCallStack(D.ShowCallStack), ImmediateDiag(D.ImmediateDiag),
+      PartialDiagId(D.PartialDiagId) {
+  // Clean the previous diagnostics.
+  D.ShowCallStack = false;
+  D.ImmediateDiag.reset();
+  D.PartialDiagId.reset();
+}
+
 Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
   if (ImmediateDiag) {
     // Emit our diagnostic and, if it was a warning or error, output a callstack
diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp
index 090d9431be..f314cb0abe 100644
--- a/lib/Sema/SemaType.cpp
+++ b/lib/Sema/SemaType.cpp
@@ -2250,15 +2250,13 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM,
   }
 
   if (T->isVariableArrayType() && !Context.getTargetInfo().isVLASupported()) {
-    if (getLangOpts().CUDA) {
-      // CUDA device code doesn't support VLAs.
-      CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget();
-    } else if (!getLangOpts().OpenMP ||
-               shouldDiagnoseTargetSupportFromOpenMP()) {
-      // Some targets don't support VLAs.
-      Diag(Loc, diag::err_vla_unsupported);
-      return QualType();
-    }
+    // CUDA device code and some other targets don't support VLAs.
+    targetDiag(Loc, (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
+                        ? diag::err_cuda_vla
+                        : diag::err_vla_unsupported)
+        << ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
+                ? CurrentCUDATarget()
+                : CFT_InvalidTarget);
   }
 
   // If this is not C99, extwarn about VLA's and C99 array size modifiers.
diff --git a/test/OpenMP/target_vla_messages.cpp b/test/OpenMP/target_vla_messages.cpp
index b744081e98..30a2751724 100644
--- a/test/OpenMP/target_vla_messages.cpp
+++ b/test/OpenMP/target_vla_messages.cpp
@@ -47,7 +47,7 @@ void target_template(int arg) {
 #pragma omp target
   {
 #ifdef NO_VLA
-    // expected-error@+2 {{variable length arrays are not supported for the current target}}
+    // expected-error@+2 2 {{variable length arrays are not supported for the current target}}
 #endif
     T vla[arg];
   }
@@ -73,6 +73,9 @@ void target(int arg) {
     }
   }
 
+#ifdef NO_VLA
+    // expected-note@+2 {{in instantiation of function template specialization 'target_template<long>' requested here}}
+#endif
   target_template<long>(arg);
 }
 
diff --git a/test/SemaCUDA/vla.cu b/test/SemaCUDA/vla.cu
index b65ae5e5fe..cf3054cd8e 100644
--- a/test/SemaCUDA/vla.cu
+++ b/test/SemaCUDA/vla.cu
@@ -1,5 +1,9 @@
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -verify %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -verify -DHOST %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux -verify -DHOST %s
+
+#ifndef __CUDA_ARCH__
+// expected-no-diagnostics
+#endif
 
 #include "Inputs/cuda.h"
 
@@ -8,7 +12,10 @@ void host(int n) {
 }
 
 __device__ void device(int n) {
-  int x[n];  // expected-error {{cannot use variable-length arrays in __device__ functions}}
+  int x[n];
+#ifdef __CUDA_ARCH__
+  // expected-error@-2 {{cannot use variable-length arrays in __device__ functions}}
+#endif
 }
 
 __host__ __device__ void hd(int n) {