[CUDA] Rework tests now that we emit deferred diagnostics during sema. Test-only...

Summary:
Previously we had to split out a lot of our tests into a test that
checked only immediate errors and a test that checked only deferred
errors.  This was because, if you emitted any immediate errors, we
wouldn't run codegen, where the deferred errors were emitted.

We've fixed this, and now emit deferred errors during sema.  This lets
us merge a bunch of tests, and lets us convert some other tests to
-fsyntax-only.

Reviewers: tra

Subscribers: cfe-commits

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

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

diff --git a/test/PCH/pragma-cuda-force-host-device.cu b/test/PCH/pragma-cuda-force-host-device.cu
index dc006be9609fcee4cff0b7b0c53bc80c9abcef35..1eaa453ad78a94f877e2e238b23220bca54d3f72 100644 (file)
--- a/test/PCH/pragma-cuda-force-host-device.cu
+++ b/test/PCH/pragma-cuda-force-host-device.cu
@@ -1,5 +1,5 @@
 // RUN: %clang_cc1 -emit-pch %s -o %t
-// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -S -o /dev/null %s
+// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -fsyntax-only %s
 
 #ifndef HEADER
 #define HEADER

diff --git a/test/Parser/cuda-force-host-device-templates.cu b/test/Parser/cuda-force-host-device-templates.cu
index 68ec9c8e67345a5336665dc919fcc6c482d4036e..315de1cf67a1b8e508aba553dc8ac00eecb4deaf 100644 (file)
--- a/test/Parser/cuda-force-host-device-templates.cu
+++ b/test/Parser/cuda-force-host-device-templates.cu
@@ -1,8 +1,7 @@
-// RUN: %clang_cc1 -std=c++14 -S -verify -fcuda-is-device %s -o /dev/null
+// RUN: %clang_cc1 -std=c++14 -fsyntax-only -verify -fcuda-is-device %s
 
 // Check how the force_cuda_host_device pragma interacts with template
-// instantiations.  The errors here are emitted at codegen, so we can't do
-// -fsyntax-only.
+// instantiations.
 
 template <typename T>
 auto foo() {  // expected-note {{declared here}}

diff --git a/test/SemaCUDA/device-var-init.cu b/test/SemaCUDA/device-var-init.cu
index 122dfca4232907c313e358815b468412f20df192..71f2352843b201f83ee7b40c9936dcd57c05502d 100644 (file)
--- a/test/SemaCUDA/device-var-init.cu
+++ b/test/SemaCUDA/device-var-init.cu
@@ -213,3 +213,15 @@ __device__ void df_sema() {
   static int v;
   // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
 }
+
+__host__ __device__ void hd_sema() {
+  static int x = 42;
+#ifdef __CUDA_ARCH__
+  // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+#endif
+}
+
+inline __host__ __device__ void hd_emitted_host_only() {
+  static int x = 42; // no error on device because this is never codegen'ed there.
+}
+void call_hd_emitted_host_only() { hd_emitted_host_only(); }

diff --git a/test/SemaCUDA/exceptions-host-device.cu b/test/SemaCUDA/exceptions-host-device.cu
deleted file mode 100644 (file)
index 9e18634..0000000
--- a/test/SemaCUDA/exceptions-host-device.cu
+++ /dev/null
@@ -1,38 +0,0 @@
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null
-// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null
-
-#include "Inputs/cuda.h"
-
-// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
-// function if and only if it's codegen'ed for device.
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void hd1() {
-  throw NULL;
-  try {} catch(void*) {}
-#ifndef HOST
-  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
-  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
-#endif
-}
-
-// No error, never instantiated on device.
-inline __host__ __device__ void hd2() {
-  throw NULL;
-  try {} catch(void*) {}
-}
-void call_hd2() { hd2(); }
-
-// Error, instantiated on device.
-inline __host__ __device__ void hd3() {
-  throw NULL;
-  try {} catch(void*) {}
-#ifndef HOST
-  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
-  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
-#endif
-}
-__device__ void call_hd3() { hd3(); }

diff --git a/test/SemaCUDA/exceptions.cu b/test/SemaCUDA/exceptions.cu
index 9ed9b6977202ca18761c0281e9954716430319ea..73d2b9d084e6fa27d67fae469538c9006f1275f9 100644 (file)
--- a/test/SemaCUDA/exceptions.cu
+++ b/test/SemaCUDA/exceptions.cu
@@ -19,3 +19,34 @@ __global__ void kernel() {
   try {} catch(void*) {}
   // expected-error@-1 {{cannot use 'try' in __global__ function}}
 }
+
+// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
+// function if and only if it's codegen'ed for device.
+
+__host__ __device__ void hd1() {
+  throw NULL;
+  try {} catch(void*) {}
+#ifdef __CUDA_ARCH__
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
+#endif
+}
+
+// No error, never instantiated on device.
+inline __host__ __device__ void hd2() {
+  throw NULL;
+  try {} catch(void*) {}
+}
+void call_hd2() { hd2(); }
+
+// Error, instantiated on device.
+inline __host__ __device__ void hd3() {
+  throw NULL;
+  try {} catch(void*) {}
+#ifdef __CUDA_ARCH__
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
+#endif
+}
+
+__device__ void call_hd3() { hd3(); }

diff --git a/test/SemaCUDA/function-overload-hd.cu b/test/SemaCUDA/function-overload-hd.cu
deleted file mode 100644 (file)
index 1b9beab..0000000
--- a/test/SemaCUDA/function-overload-hd.cu
+++ /dev/null
@@ -1,31 +0,0 @@
-// REQUIRES: x86-registered-target
-// REQUIRES: nvptx-registered-target
-
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -S -o /dev/null -verify \
-// RUN:   -verify-ignore-unexpected=note %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -S -o /dev/null -fcuda-is-device \
-// RUN:   -verify -verify-ignore-unexpected=note %s
-
-#include "Inputs/cuda.h"
-
-// FIXME: Merge into function-overload.cu once deferred errors can be emitted
-// when non-deferred errors are present.
-
-#if !defined(__CUDA_ARCH__)
-//expected-no-diagnostics
-#endif
-
-typedef void (*GlobalFnPtr)();  // __global__ functions must return void.
-
-__global__ void g() {}
-
-__host__ __device__ void hd() {
-  GlobalFnPtr fp_g = g;
-#if defined(__CUDA_ARCH__)
-  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif
-  g<<<0,0>>>();
-#if defined(__CUDA_ARCH__)
-  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
-#endif  // __CUDA_ARCH__
-}

diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
index 4545a808759279f34d48629919ad69946eae4bff..161a54ef35fcd47d2c725d9399b90b858e94cdc8 100644 (file)
--- a/test/SemaCUDA/function-overload.cu
+++ b/test/SemaCUDA/function-overload.cu
@@ -193,12 +193,22 @@ __host__ __device__ void hostdevicef() {
   CurrentFnPtr fp_cdh = cdh;
   CurrentReturnTy ret_cdh = cdh();
 
+  GlobalFnPtr fp_g = g;
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
+
   g();
 #if defined (__CUDA_ARCH__)
   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 #else
   // expected-error@-4 {{call to global function g not configured}}
 #endif
+
+  g<<<0,0>>>();
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+#endif
 }
 
 // Test for address of overloaded function resolution in the global context.

diff --git a/test/SemaCUDA/implicit-device-lambda-hd.cu b/test/SemaCUDA/implicit-device-lambda-hd.cu
deleted file mode 100644 (file)
index 6cd0e96..0000000
--- a/test/SemaCUDA/implicit-device-lambda-hd.cu
+++ /dev/null
@@ -1,27 +0,0 @@
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -verify-ignore-unexpected=note \
-// RUN:   -S -o /dev/null %s
-// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=note \
-// RUN:   -DHOST -S -o /dev/null %s
-#include "Inputs/cuda.h"
-
-__host__ __device__ void hd_fn() {
-  auto f1 = [&] {};
-  f1(); // implicitly __host__ __device__
-
-  auto f2 = [&] __device__ {};
-  f2();
-#ifdef HOST
-  // expected-error@-2 {{reference to __device__ function}}
-#endif
-
-  auto f3 = [&] __host__ {};
-  f3();
-#ifndef HOST
-  // expected-error@-2 {{reference to __host__ function}}
-#endif
-
-  auto f4 = [&] __host__ __device__ {};
-  f4();
-}
-
-

diff --git a/test/SemaCUDA/implicit-device-lambda.cu b/test/SemaCUDA/implicit-device-lambda.cu
index be1babe82293de2786b62f18a9b72e5143a74e66..8e5b7ddddb8f6edf5e7bf4fddf589c21436726e7 100644 (file)
--- a/test/SemaCUDA/implicit-device-lambda.cu
+++ b/test/SemaCUDA/implicit-device-lambda.cu
@@ -76,6 +76,26 @@ __host__ void host_fn() {
   f4();
 }
 
+__host__ __device__ void hd_fn() {
+  auto f1 = [&] {};
+  f1(); // implicitly __host__ __device__
+
+  auto f2 = [&] __device__ {};
+  f2();
+#ifndef __CUDA_ARCH__
+  // expected-error@-2 {{reference to __device__ function}}
+#endif
+
+  auto f3 = [&] __host__ {};
+  f3();
+#ifdef __CUDA_ARCH__
+  // expected-error@-2 {{reference to __host__ function}}
+#endif
+
+  auto f4 = [&] __host__ __device__ {};
+  f4();
+}
+
 // The special treatment above only applies to lambdas.
 __device__ void foo() {
   struct X {

diff --git a/test/SemaCUDA/static-vars-hd.cu b/test/SemaCUDA/static-vars-hd.cu
deleted file mode 100644 (file)
index 70cc041..0000000
--- a/test/SemaCUDA/static-vars-hd.cu
+++ /dev/null
@@ -1,20 +0,0 @@
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s
-// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s
-
-#include "Inputs/cuda.h"
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void f() {
-  static int x = 42;
-#ifndef HOST
-  // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
-#endif
-}
-
-inline __host__ __device__ void g() {
-  static int x = 42; // no error on device because this is never codegen'ed there.
-}
-void call_g() { g(); }

diff --git a/test/SemaCUDA/vla-host-device.cu b/test/SemaCUDA/vla-host-device.cu
deleted file mode 100644 (file)
index 0f54bdf..0000000
--- a/test/SemaCUDA/vla-host-device.cu
+++ /dev/null
@@ -1,21 +0,0 @@
-// RUN: %clang_cc1 -fcuda-is-device -verify -S %s -o /dev/null
-// RUN: %clang_cc1 -verify -DHOST %s -S -o /dev/null
-
-#include "Inputs/cuda.h"
-
-#ifdef HOST
-// expected-no-diagnostics
-#endif
-
-__host__ __device__ void hd(int n) {
-  int x[n];
-#ifndef HOST
-  // expected-error@-2 {{cannot use variable-length arrays in __host__ __device__ functions}}
-#endif
-}
-
-// No error because never codegen'ed for device.
-__host__ __device__ inline void hd_inline(int n) {
-  int x[n];
-}
-void call_hd_inline() { hd_inline(42); }

diff --git a/test/SemaCUDA/vla.cu b/test/SemaCUDA/vla.cu
index 7c73d9d91bf651412ee83c2c6646b95374f1ae79..f0d1ba595d935d38b9d516c8ae5a6a7da1da50b5 100644 (file)
--- a/test/SemaCUDA/vla.cu
+++ b/test/SemaCUDA/vla.cu
@@ -10,3 +10,16 @@ void host(int n) {
 __device__ void device(int n) {
   int x[n];  // expected-error {{cannot use variable-length arrays in __device__ functions}}
 }
+
+__host__ __device__ void hd(int n) {
+  int x[n];
+#ifdef __CUDA_ARCH__
+  // expected-error@-2 {{cannot use variable-length arrays in __host__ __device__ functions}}
+#endif
+}
+
+// No error because never codegen'ed for device.
+__host__ __device__ inline void hd_inline(int n) {
+  int x[n];
+}
+void call_hd_inline() { hd_inline(42); }