From: Peter Collingbourne Date: Sun, 2 Oct 2011 23:49:40 +0000 (+0000) Subject: CUDA: diagnose invalid calls across targets X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=78dd67e78c50a7abdc7c358e5eac1770d5fea22a;p=clang CUDA: diagnose invalid calls across targets git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@140978 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index 7c6d23902a..8d4ea8d610 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -1920,6 +1920,17 @@ def note_ovl_candidate_bad_base_to_derived_conv : Note<"candidate " "%select{base class pointer|superclass|base class object of type}2 %3 to " "%select{derived class pointer|subclass|derived class reference}2 %4 for " "%ordinal5 argument">; +def note_ovl_candidate_bad_target : Note< + "candidate %select{function|function|constructor|" + "function |function |constructor |" + "constructor (the implicit default constructor)|" + "constructor (the implicit copy constructor)|" + "constructor (the implicit move constructor)|" + "function (the implicit copy assignment operator)|" + "function (the implicit move assignment operator)|" + "constructor (inherited)}0 not viable: call to " + "%select{__device__|__global__|__host__|__host__ __device__}1 function from" + " %select{__device__|__global__|__host__|__host__ __device__}2 function">; def note_ambiguous_type_conversion: Note< "because of ambiguity in conversion of %0 to %1">; @@ -3992,6 +4003,9 @@ def err_kern_call_not_global_function : Error< "kernel call to non-global function %0">; def err_global_call_not_config : Error< "call to global function %0 not configured">; +def err_ref_bad_target : Error< + "reference to %select{__device__|__global__|__host__|__host__ __device__}0 " + "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">; def err_cannot_pass_objc_interface_to_vararg : Error< diff --git a/include/clang/Sema/Overload.h b/include/clang/Sema/Overload.h index e4f923703b..2dd85d5b2a 100644 --- a/include/clang/Sema/Overload.h +++ b/include/clang/Sema/Overload.h @@ -527,7 +527,12 @@ namespace clang { /// This conversion function template specialization candidate is not /// viable because the final conversion was not an exact match. - ovl_fail_final_conversion_not_exact + ovl_fail_final_conversion_not_exact, + + /// (CUDA) This candidate was not viable because the callee + /// was not accessible from the caller's target (i.e. host->device, + /// global->host, device->host). + ovl_fail_bad_target }; /// OverloadCandidate - A single candidate in an overload set (C++ 13.3). diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h index bf54bdc575..9073604d5f 100644 --- a/include/clang/Sema/Sema.h +++ b/include/clang/Sema/Sema.h @@ -5878,6 +5878,23 @@ public: QualType FieldTy, const Expr *BitWidth, bool *ZeroWidth = 0); + enum CUDAFunctionTarget { + CFT_Device, + CFT_Global, + CFT_Host, + CFT_HostDevice + }; + + CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D); + + bool CheckCUDATarget(CUDAFunctionTarget CallerTarget, + CUDAFunctionTarget CalleeTarget); + + bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee) { + return CheckCUDATarget(IdentifyCUDATarget(Caller), + IdentifyCUDATarget(Callee)); + } + /// \name Code completion //@{ /// \brief Describes the context in which code completion occurs. diff --git a/lib/Sema/SemaDeclCXX.cpp b/lib/Sema/SemaDeclCXX.cpp index 73fd18889e..bc0383165d 100644 --- a/lib/Sema/SemaDeclCXX.cpp +++ b/lib/Sema/SemaDeclCXX.cpp @@ -10883,3 +10883,44 @@ void Sema::CheckDelegatingCtorCycles() { for (CI = Invalid.begin(), CE = Invalid.end(); CI != CE; ++CI) (*CI)->setInvalidDecl(); } + +/// IdentifyCUDATarget - Determine the CUDA compilation target for this function +Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { + // Implicitly declared functions (e.g. copy constructors) are + // __host__ __device__ + if (D->isImplicit()) + return CFT_HostDevice; + + if (D->hasAttr()) + return CFT_Global; + + if (D->hasAttr()) { + if (D->hasAttr()) + return CFT_HostDevice; + else + return CFT_Device; + } + + return CFT_Host; +} + +bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, + CUDAFunctionTarget CalleeTarget) { + // CUDA B.1.1 "The __device__ qualifier declares a function that is... + // Callable from the device only." + if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device) + return true; + + // CUDA B.1.2 "The __global__ qualifier declares a function that is... + // Callable from the host only." + // CUDA B.1.3 "The __host__ qualifier declares a function that is... + // Callable from the host only." + if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) && + (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)) + return true; + + if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) + return true; + + return false; +} diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp index 46d9fe64e0..824e82b9d3 100644 --- a/lib/Sema/SemaExpr.cpp +++ b/lib/Sema/SemaExpr.cpp @@ -1379,6 +1379,20 @@ ExprResult Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK, const DeclarationNameInfo &NameInfo, const CXXScopeSpec *SS) { + if (getLangOptions().CUDA) + if (const FunctionDecl *Caller = dyn_cast(CurContext)) + if (const FunctionDecl *Callee = dyn_cast(D)) { + CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller), + CalleeTarget = IdentifyCUDATarget(Callee); + if (CheckCUDATarget(CallerTarget, CalleeTarget)) { + Diag(NameInfo.getLoc(), diag::err_ref_bad_target) + << CalleeTarget << D->getIdentifier() << CallerTarget; + Diag(D->getLocation(), diag::note_previous_decl) + << D->getIdentifier(); + return ExprError(); + } + } + MarkDeclarationReferenced(NameInfo.getLoc(), D); Expr *E = DeclRefExpr::Create(Context, diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp index 0c9083ef2f..836548aea7 100644 --- a/lib/Sema/SemaOverload.cpp +++ b/lib/Sema/SemaOverload.cpp @@ -4220,6 +4220,15 @@ Sema::AddOverloadCandidate(FunctionDecl *Function, return; } + // (CUDA B.1): Check for invalid calls between targets. + if (getLangOptions().CUDA) + if (const FunctionDecl *Caller = dyn_cast(CurContext)) + if (CheckCUDATarget(Caller, Function)) { + Candidate.Viable = false; + Candidate.FailureKind = ovl_fail_bad_target; + return; + } + // Determine the implicit conversion sequences for each of the // arguments. Candidate.Conversions.resize(NumArgs); @@ -7189,6 +7198,21 @@ void DiagnoseBadDeduction(Sema &S, OverloadCandidate *Cand, } } +/// CUDA: diagnose an invalid call across targets. +void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) { + FunctionDecl *Caller = cast(S.CurContext); + FunctionDecl *Callee = Cand->Function; + + Sema::CUDAFunctionTarget CallerTarget = S.IdentifyCUDATarget(Caller), + CalleeTarget = S.IdentifyCUDATarget(Callee); + + std::string FnDesc; + OverloadCandidateKind FnKind = ClassifyOverloadCandidate(S, Callee, FnDesc); + + S.Diag(Callee->getLocation(), diag::note_ovl_candidate_bad_target) + << (unsigned) FnKind << CalleeTarget << CallerTarget; +} + /// Generates a 'note' diagnostic for an overload candidate. We've /// already generated a primary error at the call site. /// @@ -7248,6 +7272,9 @@ void NoteFunctionCandidate(Sema &S, OverloadCandidate *Cand, // those conditions and diagnose them well. return S.NoteOverloadCandidate(Fn); } + + case ovl_fail_bad_target: + return DiagnoseBadTarget(S, Cand); } } @@ -7780,6 +7807,11 @@ private: return false; if (FunctionDecl *FunDecl = dyn_cast(Fn)) { + if (S.getLangOptions().CUDA) + if (FunctionDecl *Caller = dyn_cast(S.CurContext)) + if (S.CheckCUDATarget(Caller, FunDecl)) + return false; + QualType ResultTy; if (Context.hasSameUnqualifiedType(TargetFunctionType, FunDecl->getType()) || diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h index e3aeb99ed2..26a8df0440 100644 --- a/test/SemaCUDA/cuda.h +++ b/test/SemaCUDA/cuda.h @@ -10,7 +10,7 @@ struct dim3 { unsigned x, y, z; - dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {} + __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {} }; typedef struct cudaStream *cudaStream_t; diff --git a/test/SemaCUDA/function-target.cu b/test/SemaCUDA/function-target.cu new file mode 100644 index 0000000000..c7a55e2fad --- /dev/null +++ b/test/SemaCUDA/function-target.cu @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +#include "cuda.h" + +__host__ void h1h(void); +__device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}} +__host__ __device__ void h1hd(void); +__global__ void h1g(void); + +struct h1ds { // expected-note {{requires 1 argument}} + __device__ h1ds(); // expected-note {{candidate constructor not viable: call to __device__ function from __host__ function}} +}; + +__host__ void h1(void) { + h1h(); + h1d(); // expected-error {{no matching function}} + h1hd(); + h1g<<<1, 1>>>(); + h1ds x; // expected-error {{no matching constructor}} +} + +__host__ void d1h(void); // expected-note {{candidate function not viable: call to __host__ function from __device__ function}} +__device__ void d1d(void); +__host__ __device__ void d1hd(void); +__global__ void d1g(void); // expected-note {{'d1g' declared here}} + +__device__ void d1(void) { + d1h(); // expected-error {{no matching function}} + d1d(); + d1hd(); + d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}} +} + +__host__ void hd1h(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}} +__device__ void hd1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}} +__host__ __device__ void hd1hd(void); +__global__ void hd1g(void); // expected-note {{'hd1g' declared here}} + +__host__ __device__ void hd1(void) { + hd1h(); // expected-error {{no matching function}} + hd1d(); // expected-error {{no matching function}} + hd1hd(); + hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}} +}