12 files changed, 89 insertions, 141 deletions
diff --git a/clang/test/PCH/pragma-cuda-force-host-device.cu b/clang/test/PCH/pragma-cuda-force-host-device.cu
index dc006be9609..1eaa453ad78 100644
--- a/clang/test/PCH/pragma-cuda-force-host-device.cu
+++ b/clang/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/clang/test/Parser/cuda-force-host-device-templates.cu b/clang/test/Parser/cuda-force-host-device-templates.cu
index 68ec9c8e673..315de1cf67a 100644
--- a/clang/test/Parser/cuda-force-host-device-templates.cu
+++ b/clang/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/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu
index 122dfca4232..71f2352843b 100644
--- a/clang/test/SemaCUDA/device-var-init.cu
+++ b/clang/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/clang/test/SemaCUDA/exceptions-host-device.cu b/clang/test/SemaCUDA/exceptions-host-device.cu
deleted file mode 100644
index 9e18634e905..00000000000
--- a/clang/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/clang/test/SemaCUDA/exceptions.cu b/clang/test/SemaCUDA/exceptions.cu
index 9ed9b697720..73d2b9d084e 100644
--- a/clang/test/SemaCUDA/exceptions.cu
+++ b/clang/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/clang/test/SemaCUDA/function-overload-hd.cu b/clang/test/SemaCUDA/function-overload-hd.cu
deleted file mode 100644
index 1b9beab7268..00000000000
--- a/clang/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/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu
index 4545a808759..161a54ef35f 100644
--- a/clang/test/SemaCUDA/function-overload.cu
+++ b/clang/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/clang/test/SemaCUDA/implicit-device-lambda-hd.cu b/clang/test/SemaCUDA/implicit-device-lambda-hd.cu
deleted file mode 100644
index 6cd0e96af8c..00000000000
--- a/clang/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/clang/test/SemaCUDA/implicit-device-lambda.cu b/clang/test/SemaCUDA/implicit-device-lambda.cu
index be1babe8229..8e5b7ddddb8 100644
--- a/clang/test/SemaCUDA/implicit-device-lambda.cu
+++ b/clang/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/clang/test/SemaCUDA/static-vars-hd.cu b/clang/test/SemaCUDA/static-vars-hd.cu
deleted file mode 100644
index 70cc0417f0e..00000000000
--- a/clang/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/clang/test/SemaCUDA/vla-host-device.cu b/clang/test/SemaCUDA/vla-host-device.cu
deleted file mode 100644
index 0f54bdf4917..00000000000
--- a/clang/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/clang/test/SemaCUDA/vla.cu b/clang/test/SemaCUDA/vla.cu
index 7c73d9d91bf..f0d1ba595d9 100644
--- a/clang/test/SemaCUDA/vla.cu
+++ b/clang/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); }