Fix for PR20886 - enforce CUDA target match in method calls

[eliben]

Friendly ping

[eliben]

Friendly ping!

[pcc]

test/SemaCUDA/method-target.cu

22

Can we check for a more specific error message here?

[eliben]

Addressing Peter's review

[eliben]

Done, PTAL

test/SemaCUDA/method-target.cu

22

Done

[pcc]

LGTM

[eliben]

Thanks for the review, Peter. Committed in Clang r218482.

[eliben]

Fixed

lib/Sema/SemaOverload.cpp

                                         Method, ActingContext);
     if (Candidate.Conversions[0].isBad()) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_bad_conversion;
       return;
     }
   }
 
+  // (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) {
     if (ArgIdx < NumParams) {
       // (C++ 13.3.2p3): for F to be a viable function, there shall
       // exist for each argument an implicit conversion sequence
       // (13.3.3.1) that converts that argument to the corresponding
       // parameter of F.

   ExprValueKind VK = Expr::getValueKindForType(ResultType);
   ResultType = ResultType.getNonLValueExprType(Context);
 
   assert(Method && "Member call to something that isn't a method?");
   CXXMemberCallExpr *TheCall =
     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))
     return ExprError();
 
   // Convert the object argument (for a non-static member function call).
   // We only need to do this if there was actually an overload; otherwise
   // it was done at lookup.

test/SemaCUDA/method-target.cu

+// 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}}
+}