[CUDA] Remove three obsolete CUDA cc1 flags.

Summary:
* -fcuda-target-overloads

  Previously unconditionally set to true by the driver.  Necessary for
  correct functioning of the compiler -- our CUDA headers wrapper won't
  compile without this.

* -fcuda-disable-target-call-checks

  Previously unconditionally set to true by the driver.  Necessary to
  compile almost any external CUDA code -- almost all libraries assume
  that host+device code can call host or device functions.

* -fcuda-allow-host-calls-from-host-device

  No effect when target overloading is enabled.

Reviewers: tra

Subscribers: rsmith, cfe-commits

Differential Revision: http://reviews.llvm.org/D18416

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

7 files changed, 9 insertions, 140 deletions

diff --git a/test/SemaCUDA/builtins.cu b/test/SemaCUDA/builtins.cu
index 32b575862c..7e6d014c3f 100644
--- a/test/SemaCUDA/builtins.cu
+++ b/test/SemaCUDA/builtins.cu
@@ -7,10 +7,10 @@
 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple x86_64-unknown-unknown \
 // RUN:     -aux-triple nvptx64-unknown-cuda \
-// RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN:     -fsyntax-only -verify %s
 // RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
 // RUN:     -aux-triple x86_64-unknown-unknown \
-// RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
+// RUN:     -fsyntax-only -verify %s
 
 #if !(defined(__amd64__) && defined(__PTX__))
 #error "Expected to see preprocessor macros from both sides of compilation."
diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
index 8de9e63bd6..3c78600b17 100644
--- a/test/SemaCUDA/function-overload.cu
+++ b/test/SemaCUDA/function-overload.cu
@@ -1,18 +1,8 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// Make sure we handle target overloads correctly.
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
-// RUN:    -fsyntax-only -fcuda-target-overloads -verify %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
-// RUN:    -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s
-
-// Check target overloads handling with disabled call target checks.
-// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
-// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
-// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
-// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
-// RUN:    -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
@@ -180,23 +170,11 @@ __host__ __device__ void hostdevicef() {
   DeviceReturnTy ret_d = d();
   DeviceFnPtr fp_cd = cd;
   DeviceReturnTy ret_cd = cd();
-#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
-  // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
-  // expected-error@-5 {{no matching function for call to 'd'}}
-  // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
-  // expected-error@-5 {{no matching function for call to 'cd'}}
-#endif
 
   HostFnPtr fp_h = h;
   HostReturnTy ret_h = h();
   HostFnPtr fp_ch = ch;
   HostReturnTy ret_ch = ch();
-#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
-  // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
-  // expected-error@-5 {{no matching function for call to 'h'}}
-  // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
-  // expected-error@-5 {{no matching function for call to 'ch'}}
-#endif
 
   CurrentFnPtr fp_dh = dh;
   CurrentReturnTy ret_dh = dh();
@@ -372,13 +350,7 @@ __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {
 
 __host__ __device__ void test_host_device_calls_hd_template() {
   HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
-
-#if defined(__CUDA_ARCH__) && !defined(NOCHECKS)
-  typedef HostDeviceReturnTy ExpectedReturnTy;
-#else
-  typedef TemplateReturnTy ExpectedReturnTy;
-#endif
-  ExpectedReturnTy ret2 = template_vs_hd_function(1);
+  TemplateReturnTy ret2 = template_vs_hd_function(1);
 }
 
 __host__ void test_host_calls_hd_template() {
@@ -401,11 +373,9 @@ __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturn
 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }
 
-__host__ __device__ void test_host_device_nochecks_overloading() {
-#ifdef NOCHECKS
+__host__ __device__ void test_host_device_single_side_overloading() {
   DeviceReturnTy ret1 = device_only_function(1);
   DeviceReturnTy2 ret2 = device_only_function(1.0f);
   HostReturnTy ret3 = host_only_function(1);
   HostReturnTy2 ret4 = host_only_function(1.0f);
-#endif
 }
diff --git a/test/SemaCUDA/function-target-disabled-check.cu b/test/SemaCUDA/function-target-disabled-check.cu
deleted file mode 100644
index 979d4edbf8..0000000000
--- a/test/SemaCUDA/function-target-disabled-check.cu
+++ /dev/null
@@ -1,26 +0,0 @@
-// Test that we can disable cross-target call checks in Sema with the
-// -fcuda-disable-target-call-checks flag. Without this flag we'd get a bunch
-// of errors here, since there are invalid cross-target calls present.
-
-// RUN: %clang_cc1 -fsyntax-only -verify %s -fcuda-disable-target-call-checks
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -fcuda-disable-target-call-checks
-
-// expected-no-diagnostics
-
-#define __device__ __attribute__((device))
-#define __global__ __attribute__((global))
-#define __host__ __attribute__((host))
-
-__attribute__((host)) void h1();
-
-__attribute__((device)) void d1() {
-  h1();
-}
-
-__attribute__((host)) void h2() {
-  d1();
-}
-
-__attribute__((global)) void g1() {
-  h2();
-}
diff --git a/test/SemaCUDA/function-target-hd.cu b/test/SemaCUDA/function-target-hd.cu
deleted file mode 100644
index 685f4f9cda..0000000000
--- a/test/SemaCUDA/function-target-hd.cu
+++ /dev/null
@@ -1,71 +0,0 @@
-// Test the Sema analysis of caller-callee relationships of host device
-// functions when compiling CUDA code. There are 4 permutations of this test as
-// host and device compilation are separate compilation passes, and clang has
-// an option to allow host calls from host device functions. __CUDA_ARCH__ is
-// defined when compiling for the device and TEST_WARN_HD when host calls are
-// allowed from host device functions. So for example, if __CUDA_ARCH__ is
-// defined and TEST_WARN_HD is not then device compilation is happening but
-// host device functions are not allowed to call device functions.
-
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
-
-#include "Inputs/cuda.h"
-
-__host__ void hd1h(void);
-#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
-// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-#endif
-__device__ void hd1d(void);
-#ifndef __CUDA_ARCH__
-// expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-#endif
-__host__ void hd1hg(void);
-__device__ void hd1dg(void);
-#ifdef __CUDA_ARCH__
-__host__ void hd1hig(void);
-#if !defined(TEST_WARN_HD)
-// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-#endif
-#else
-__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-#endif
-__host__ __device__ void hd1hd(void);
-__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
-
-__host__ __device__ void hd1(void) {
-#if defined(TEST_WARN_HD) && defined(__CUDA_ARCH__)
-// expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
-// expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
-#endif
-  hd1d();
-#ifndef __CUDA_ARCH__
-// expected-error@-2 {{no matching function}}
-#endif
-  hd1h();
-#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
-// expected-error@-2 {{no matching function}}
-#endif
-
-  // No errors as guarded
-#ifdef __CUDA_ARCH__
-  hd1d();
-#else
-  hd1h();
-#endif
-
-  // Errors as incorrectly guarded
-#ifndef __CUDA_ARCH__
-  hd1dig(); // expected-error {{no matching function}}
-#else
-  hd1hig();
-#ifndef TEST_WARN_HD
-// expected-error@-2 {{no matching function}}
-#endif
-#endif
-
-  hd1hd();
-  hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
-}
diff --git a/test/SemaCUDA/implicit-intrinsic.cu b/test/SemaCUDA/implicit-intrinsic.cu
index 0793d64b10..dba26c5539 100644
--- a/test/SemaCUDA/implicit-intrinsic.cu
+++ b/test/SemaCUDA/implicit-intrinsic.cu
@@ -1,7 +1,5 @@
 // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
 // RUN:     -fsyntax-only -verify %s
-// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \
-// RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
 
 #include "Inputs/cuda.h"
 
diff --git a/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
index f038c376ff..7aa1dd3f20 100644
--- a/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
+++ b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -74,13 +74,11 @@ struct B4_with_device_copy_ctor {
 struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
 };
 
-// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
-// expected-note@-4 {{implicit copy constructor inferred target collision}}
-// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-3 {{copy constructor of 'C4_with_collision' is implicitly deleted because base class 'B4_with_device_copy_ctor' has no copy constructor}}
 
 void hostfoo4() {
   C4_with_collision c;
-  C4_with_collision c2 = c; // expected-error {{no matching constructor}}
+  C4_with_collision c2 = c; // expected-error {{call to implicitly-deleted copy constructor of 'C4_with_collision'}}
 }
 
 //------------------------------------------------------------------------------
diff --git a/test/SemaCUDA/method-target.cu b/test/SemaCUDA/method-target.cu
index 4fa290719c..3fabfc359b 100644
--- a/test/SemaCUDA/method-target.cu
+++ b/test/SemaCUDA/method-target.cu
@@ -44,7 +44,7 @@ struct S4 {
 };
 
 __host__ __device__ void foo4(S4& s) {
-  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+  s.method();
 }
 
 //------------------------------------------------------------------------------