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 7c6d239..8d4ea8d 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -1920,6 +1920,17 @@
"%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 @@
"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 e4f9237..2dd85d5 100644
--- a/include/clang/Sema/Overload.h
+++ b/include/clang/Sema/Overload.h
@@ -527,7 +527,12 @@
/// 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 bf54bdc..9073604 100644
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -5878,6 +5878,23 @@
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 73fd188..bc03831 100644
--- a/lib/Sema/SemaDeclCXX.cpp
+++ b/lib/Sema/SemaDeclCXX.cpp
@@ -10883,3 +10883,44 @@
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<CUDAGlobalAttr>())
+ return CFT_Global;
+
+ if (D->hasAttr<CUDADeviceAttr>()) {
+ if (D->hasAttr<CUDAHostAttr>())
+ 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 46d9fe6..824e82b 100644
--- a/lib/Sema/SemaExpr.cpp
+++ b/lib/Sema/SemaExpr.cpp
@@ -1379,6 +1379,20 @@
Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK,
const DeclarationNameInfo &NameInfo,
const CXXScopeSpec *SS) {
+ if (getLangOptions().CUDA)
+ if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
+ if (const FunctionDecl *Callee = dyn_cast<FunctionDecl>(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 0c9083e..836548a 100644
--- a/lib/Sema/SemaOverload.cpp
+++ b/lib/Sema/SemaOverload.cpp
@@ -4220,6 +4220,15 @@
return;
}
+ // (CUDA B.1): Check for invalid calls between targets.
+ if (getLangOptions().CUDA)
+ if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(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 @@
}
}
+/// CUDA: diagnose an invalid call across targets.
+void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
+ FunctionDecl *Caller = cast<FunctionDecl>(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 @@
// those conditions and diagnose them well.
return S.NoteOverloadCandidate(Fn);
}
+
+ case ovl_fail_bad_target:
+ return DiagnoseBadTarget(S, Cand);
}
}
@@ -7780,6 +7807,11 @@
return false;
if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
+ if (S.getLangOptions().CUDA)
+ if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(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 e3aeb99..26a8df0 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 0000000..c7a55e2
--- /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}}
+}