r218482 - Fix PR20886 - enforce CUDA target match in method calls

[Eli Bendersky]

Author: eliben
Date: Thu Sep 25 18:59:08 2014
New Revision: 218482

URL: http://llvm.org/viewvc/llvm-project?rev=218482&view=rev
Log:
Fix PR20886 - enforce CUDA target match in method calls

http://reviews.llvm.org/D5298


Added:
    cfe/trunk/test/SemaCUDA/method-target.cu
Modified:
    cfe/trunk/lib/Sema/SemaOverload.cpp

Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=218482&r1=218481&r2=218482&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Thu Sep 25 18:59:08 2014
@@ -5977,6 +5977,15 @@ Sema::AddMethodCandidate(CXXMethodDecl *
     }
   }
 
+  // (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) {
@@ -11598,6 +11607,18 @@ Sema::BuildCallToMemberFunction(Scope *S
     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))

Added: cfe/trunk/test/SemaCUDA/method-target.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/method-target.cu?rev=218482&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/method-target.cu (added)
+++ cfe/trunk/test/SemaCUDA/method-target.cu Thu Sep 25 18:59:08 2014
@@ -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: call to __host__ function from __device__ function}}
+  void method(float) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+};
+
+__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}}
+}