[CUDA] Allow function overloads in CUDA based on host/device attributes.

The patch makes it possible to parse CUDA files that contain host/device
functions with identical signatures, but different attributes without
having to physically split source into host-only and device-only parts.

This change is needed in order to parse CUDA header files that have
a lot of name clashes with standard include files.

Gory details are in design doc here: https://goo.gl/EXnymm
Feel free to leave comments there or in this review thread.

This feature is controlled with CC1 option -fcuda-target-overloads
and is disabled by default.

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

llvm-svn: 248295

2 files changed, 531 insertions, 0 deletions

diff --git a/clang/test/CodeGenCUDA/function-overload.cu b/clang/test/CodeGenCUDA/function-overload.cu
new file mode 100644
index 00000000000..a12ef82773a
--- /dev/null
+++ b/clang/test/CodeGenCUDA/function-overload.cu
@@ -0,0 +1,214 @@
+// 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:     -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
+// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN:    -fcuda-is-device -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
+// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+typedef void (*gp_t)(void);
+
+// CHECK-HOST: @hp = global i32 ()* @_Z1hv
+// CHECK-HOST: @chp = global i32 ()* @ch
+// CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
+// CHECK-HOST: @cdhp = global i32 ()* @cdh
+// CHECK-HOST: @gp = global void ()* @_Z1gv
+
+// CHECK-BOTH-LABEL: define i32 @_Z2dhv()
+__device__ int dh(void) { return 1; }
+// CHECK-DEVICE: ret i32 1
+__host__ int dh(void) { return 2; }
+// CHECK-HOST:   ret i32 2
+
+// CHECK-BOTH-LABEL: define i32 @_Z2hdv()
+__host__ __device__ int hd(void) { return 3; }
+// CHECK-BOTH:   ret i32 3
+
+// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
+__device__ int d(void) { return 8; }
+// CHECK-DEVICE:   ret i32 8
+
+// CHECK-HOST-LABEL: define i32 @_Z1hv()
+__host__ int h(void) { return 9; }
+// CHECK-HOST:   ret i32 9
+
+// CHECK-BOTH-LABEL: define void @_Z1gv()
+__global__ void g(void) {}
+// CHECK-BOTH:   ret void
+
+// mangled names of extern "C" __host__ __device__ functions clash
+// with those of their __host__/__device__ counterparts, so
+// overloading of extern "C" functions can only happen for __host__
+// and __device__ functions -- we never codegen them in the same
+// compilation and therefore mangled name conflict is not a problem.
+
+// CHECK-BOTH-LABEL: define i32 @cdh()
+extern "C" __device__ int cdh(void) {return 10;}
+// CHECK-DEVICE:   ret i32 10
+extern "C" __host__ int cdh(void) {return 11;}
+// CHECK-HOST:     ret i32 11
+
+// CHECK-DEVICE-LABEL: define i32 @cd()
+extern "C" __device__ int cd(void) {return 12;}
+// CHECK-DEVICE:   ret i32 12
+
+// CHECK-HOST-LABEL: define i32 @ch()
+extern "C" __host__ int ch(void) {return 13;}
+// CHECK-HOST:     ret i32 13
+
+// CHECK-BOTH-LABEL: define i32 @chd()
+extern "C" __host__ __device__ int chd(void) {return 14;}
+// CHECK-BOTH:     ret i32 14
+
+// CHECK-HOST-LABEL: define void @_Z5hostfv()
+__host__ void hostf(void) {
+#if defined (NOCHECKS)
+  fp_t dp = d;   // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
+#endif
+  fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
+  fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
+  gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
+
+#if defined (NOCHECKS)
+  d();     // CHECK-HOST-NC: call i32 @_Z1dv()
+  cd();    // CHECK-HOST-NC: call i32 @cd()
+#endif
+  h();     // CHECK-HOST: call i32 @_Z1hv()
+  ch();    // CHECK-HOST: call i32 @ch()
+  dh();    // CHECK-HOST: call i32 @_Z2dhv()
+  cdh();   // CHECK-HOST: call i32 @cdh()
+  g<<<0,0>>>();  // CHECK-HOST: call void @_Z1gv()
+}
+
+// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
+__device__ void devicef(void) {
+  fp_t dp = d;   // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
+#if defined (NOCHECKS)
+  fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
+#endif
+  fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
+
+  d();     // CHECK-DEVICE: call i32 @_Z1dv()
+  cd();    // CHECK-DEVICE: call i32 @cd()
+#if defined (NOCHECKS)
+  h();     // CHECK-DEVICE-NC: call i32 @_Z1hv()
+  ch();    // CHECK-DEVICE-NC: call i32 @ch()
+#endif
+  dh();    // CHECK-DEVICE: call i32 @_Z2dhv()
+  cdh();   // CHECK-DEVICE: call i32 @cdh()
+}
+
+// CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
+__host__ __device__ void hostdevicef(void) {
+#if defined (NOCHECKS)
+  fp_t dp = d;   // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
+  fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
+#endif
+  fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
+#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
+  gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
+#endif
+
+#if defined (NOCHECKS)
+  d();     // CHECK-BOTH-NC: call i32 @_Z1dv()
+  cd();    // CHECK-BOTH-NC: call i32 @cd()
+  h();     // CHECK-BOTH-NC: call i32 @_Z1hv()
+  ch();    // CHECK-BOTH-NC: call i32 @ch()
+#endif
+  dh();    // CHECK-BOTH: call i32 @_Z2dhv()
+  cdh();   // CHECK-BOTH: call i32 @cdh()
+#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
+  g<<<0,0>>>();  // CHECK-HOST-NC: call void @_Z1gv()
+#endif
+}
+
+// Test for address of overloaded function resolution in the global context.
+fp_t hp = h;
+fp_t chp = ch;
+fp_t dhp = dh;
+fp_t cdhp = cdh;
+gp_t gp = g;
+
+int x;
+// Check constructors/destructors for D/H functions
+struct s_cd_dh {
+  __host__ s_cd_dh() { x = 11; }
+  __device__ s_cd_dh() { x = 12; }
+  __host__ ~s_cd_dh() { x = 21; }
+  __device__ ~s_cd_dh() { x = 22; }
+};
+
+struct s_cd_hd {
+  __host__ __device__ s_cd_hd() { x = 31; }
+  __host__ __device__ ~s_cd_hd() { x = 32; }
+};
+
+// CHECK-BOTH: define void @_Z7wrapperv
+#if defined(__CUDA_ARCH__)
+__device__
+#else
+__host__
+#endif
+void wrapper() {
+  s_cd_dh scddh;
+  // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
+  s_cd_hd scdhd;
+  // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
+
+  // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
+  // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
+}
+// CHECK-BOTH: ret void
+
+// Now it's time to check what's been generated for the methods we used.
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
+// CHECK-HOST:   store i32 11,
+// CHECK-DEVICE: store i32 12,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
+// CHECK-BOTH:   store i32 31,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
+// CHECK-BOTH: store i32 32,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
+// CHECK-HOST:   store i32 21,
+// CHECK-DEVICE: store i32 22,
+// CHECK-BOTH: ret void
+
diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu
new file mode 100644
index 00000000000..bd3fb508bfa
--- /dev/null
+++ b/clang/test/SemaCUDA/function-overload.cu
@@ -0,0 +1,317 @@
+// 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
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+typedef void (*gp_t)(void);
+
+// Host and unattributed functions can't be overloaded
+__host__ int hh(void) { return 1; } // expected-note {{previous definition is here}}
+int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}}
+
+// H/D overloading is OK
+__host__ int dh(void) { return 2; }
+__device__ int dh(void) { return 2; }
+
+// H/HD and D/HD are not allowed
+__host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}}
+__host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}}
+
+__host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}}
+__host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+__host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}}
+__device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}}
+
+__device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}}
+__host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Same tests for extern "C" functions
+extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}}
+extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}}
+
+// H/D overloading is OK
+extern "C" __device__ int cdh(void) {return 10;}
+extern "C" __host__ int cdh(void) {return 11;}
+
+// H/HD and D/HD overloading is not allowed.
+extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}}
+extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}}
+
+extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}}
+extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Helper functions to verify calling restrictions.
+__device__ int d(void) { return 8; }
+__host__ int h(void) { return 9; }
+__global__ void g(void) {}
+extern "C" __device__ int cd(void) {return 10;}
+extern "C" __host__ int ch(void) {return 11;}
+
+__host__ void hostf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#endif
+  fp_t hp = h;
+  fp_t chp = ch;
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+
+  d();
+  cd();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
+  // expected-error@-4 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
+#endif
+  h();
+  ch();
+  dh();
+  cdh();
+  g(); // expected-error {{call to global function g not configured}}
+  g<<<0,0>>>();
+}
+
+
+__device__ void devicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__global__ void globalf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__host__ __device__ void hostdevicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#else
+  // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+  // expected-error@-7 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#else
+  // expected-error@-9 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+  // expected-error@-10 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+#endif
+
+  dh();
+  cdh();
+  g();
+  g<<<0,0>>>();
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-3 {{call to global function g not configured}}
+#else
+  // expected-error@-5 {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
+  // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif  // __CUDA_ARCH__
+}
+
+// Test for address of overloaded function resolution in the global context.
+fp_t hp = h;
+fp_t chp = ch;
+fp_t dhp = dh;
+fp_t cdhp = cdh;
+gp_t gp = g;
+
+
+// Test overloading of destructors
+// Can't mix H and unattributed destructors
+struct d_h {
+  ~d_h() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// H/D overloading is OK
+struct d_dh {
+  __device__ ~d_dh() {}
+  __host__ ~d_dh() {}
+};
+
+// HD is OK
+struct d_hd {
+  __host__ __device__ ~d_hd() {}
+};
+
+// Mixing H/D and HD is not allowed.
+struct d_dhhd {
+  __device__ ~d_dhhd() {}
+  __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hhd {
+  __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdh {
+  __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_dhd {
+  __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdd {
+  __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
+  __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// Test overloading of member functions
+struct m_h {
+  void operator delete(void *ptr); // expected-note {{previous declaration is here}}
+  __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
+};
+
+// D/H overloading is OK
+struct m_dh {
+  __device__ void operator delete(void *ptr);
+  __host__ void operator delete(void *ptr);
+};
+
+// HD by itself is OK
+struct m_hd {
+  __device__ __host__ void operator delete(void *ptr);
+};
+
+struct m_hhd {
+  __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_hdh {
+  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_dhd {
+  __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};
+
+struct m_hdd {
+  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
+  __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+};