jlebar created this revision.
jlebar added a reviewer: tra.
jlebar added a subscriber: cfe-commits.
Specifically, this patch adds testcases for all three calls to
EraseUnwantedCUDAMatches. The addr-of-overloaded-fn test I accidentally
neutered in r264207, which moved much of
CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu.
The coverage from overloaded-delete test is new.
http://reviews.llvm.org/D21913
Files:
test/SemaCUDA/addr-of-overloaded-fn.cu
test/SemaCUDA/overloaded-delete.cu
Index: test/SemaCUDA/overloaded-delete.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/overloaded-delete.cu
@@ -0,0 +1,25 @@
+// expected-no-diagnostics
+
+// 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"
+
+struct S {
+ __host__ static void operator delete(void*, size_t) {}
+ __device__ static void operator delete(void*, size_t) {}
+};
+
+__host__ __device__ void test(S* s) {
+ // This shouldn't be ambiguous -- we call the host overload in host mode and
+ // the device overload in device mode.
+ delete s;
+}
+
+__host__ void operator delete(void *ptr) {}
+__device__ void operator delete(void *ptr) {}
+
+__host__ __device__ void test_global_delete(int *ptr) {
+ // Again, there should be no ambiguity between which operator delete we call.
+ ::delete ptr;
+}
Index: test/SemaCUDA/addr-of-overloaded-fn.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/addr-of-overloaded-fn.cu
@@ -0,0 +1,24 @@
+// expected-no-diagnostics
+
+// 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"
+
+__host__ void overload() {}
+__device__ void overload() {}
+
+__host__ __device__ void test_hd() {
+ // This should not be ambiguous -- we choose the host or the device overload
+ // depending on whether or not we're compiling for host or device.
+ void (*x)() = overload;
+}
+
+// These also shouldn't be ambiguous, but they're an easier test than the HD
+// function above.
+__host__ void test_host() {
+ void (*x)() = overload;
+}
+__device__ void test_device() {
+ void (*x)() = overload;
+}
Index: test/SemaCUDA/overloaded-delete.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/overloaded-delete.cu
@@ -0,0 +1,25 @@
+// expected-no-diagnostics
+
+// 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"
+
+struct S {
+ __host__ static void operator delete(void*, size_t) {}
+ __device__ static void operator delete(void*, size_t) {}
+};
+
+__host__ __device__ void test(S* s) {
+ // This shouldn't be ambiguous -- we call the host overload in host mode and
+ // the device overload in device mode.
+ delete s;
+}
+
+__host__ void operator delete(void *ptr) {}
+__device__ void operator delete(void *ptr) {}
+
+__host__ __device__ void test_global_delete(int *ptr) {
+ // Again, there should be no ambiguity between which operator delete we call.
+ ::delete ptr;
+}
Index: test/SemaCUDA/addr-of-overloaded-fn.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/addr-of-overloaded-fn.cu
@@ -0,0 +1,24 @@
+// expected-no-diagnostics
+
+// 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"
+
+__host__ void overload() {}
+__device__ void overload() {}
+
+__host__ __device__ void test_hd() {
+ // This should not be ambiguous -- we choose the host or the device overload
+ // depending on whether or not we're compiling for host or device.
+ void (*x)() = overload;
+}
+
+// These also shouldn't be ambiguous, but they're an easier test than the HD
+// function above.
+__host__ void test_host() {
+ void (*x)() = overload;
+}
+__device__ void test_device() {
+ void (*x)() = overload;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
