Hi pcc, rsmith,
http://reviews.llvm.org/D5298
Files:
lib/Sema/SemaOverload.cpp
test/SemaCUDA/method-target.cu
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -5977,6 +5977,15 @@
}
}
+ // (CUDA B.1): Check for invalid calls between targets.
+ if (getLangOpts().CUDA)
+ if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
+ if (CheckCUDATarget(Caller, Method)) {
+ Candidate.Viable = false;
+ Candidate.FailureKind = ovl_fail_bad_target;
+ return;
+ }
+
// Determine the implicit conversion sequences for each of the
// arguments.
for (unsigned ArgIdx = 0; ArgIdx < Args.size(); ++ArgIdx) {
@@ -11595,6 +11604,18 @@
new (Context) CXXMemberCallExpr(Context, MemExprE, Args,
ResultType, VK, RParenLoc);
+ // (CUDA B.1): Check for invalid calls between targets.
+ if (getLangOpts().CUDA) {
+ if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext)) {
+ if (CheckCUDATarget(Caller, Method)) {
+ Diag(MemExpr->getMemberLoc(), diag::err_ref_bad_target)
+ << IdentifyCUDATarget(Method) << Method->getIdentifier()
+ << IdentifyCUDATarget(Caller);
+ return ExprError();
+ }
+ }
+ }
+
// Check for a valid return type.
if (CheckCallReturnType(Method->getReturnType(), MemExpr->getMemberLoc(),
TheCall, Method))
Index: test/SemaCUDA/method-target.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/method-target.cu
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: host method called from device function
+
+struct S1 {
+ void method() {}
+};
+
+__device__ void foo1(S1& s) {
+ s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: host method called from device function, for overloaded method
+
+struct S2 {
+ void method(int) {} // expected-note {{candidate function not viable}}
+ void method(float) {} // expected-note {{candidate function not viable}}
+};
+
+__device__ void foo2(S2& s, int i, float f) {
+ s.method(f); // expected-error {{no matching member function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: device method called from host function
+
+struct S3 {
+ __device__ void method() {}
+};
+
+void foo3(S3& s) {
+ s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 4: device method called from host&device function
+
+struct S4 {
+ __device__ void method() {}
+};
+
+__host__ __device__ void foo4(S4& s) {
+ s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: overloaded operators
+
+struct S5 {
+ S5() {}
+ S5& operator=(const S5&) {return *this;} // expected-note {{candidate function not viable}}
+};
+
+__device__ void foo5(S5& s, S5& t) {
+ s = t; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: call method through pointer
+
+struct S6 {
+ void method() {}
+};
+
+__device__ void foo6(S6* s) {
+ s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
