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}} +} _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
