Author: jlebar Date: Wed Oct 19 16:03:38 2016 New Revision: 284643 URL: http://llvm.org/viewvc/llvm-project?rev=284643&view=rev Log: [CUDA] Emit errors for wrong-side calls made on the same line as non-wrong-side calls.
Summary: This fixes two related bugs: 1) Previously, if you had a non-wrong side call at some source code location L, we wouldn't emit errors for wrong-side calls that appeared at L. 2) We'd only emit one wrong-side error per source code location, when we actually want to emit it twice if we hit this line more than once due to e.g. template instantiation. Reviewers: tra Subscribers: rnk, cfe-commits Differential Revision: https://reviews.llvm.org/D25702 Added: cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu Modified: cfe/trunk/include/clang/Sema/Sema.h cfe/trunk/lib/Sema/SemaCUDA.cpp Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=284643&r1=284642&r2=284643&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Wed Oct 19 16:03:38 2016 @@ -9252,10 +9252,10 @@ public: llvm::DenseMap<const FunctionDecl *, std::vector<PartialDiagnosticAt>> CUDADeferredDiags; - /// Raw encodings of SourceLocations for which CheckCUDACall has emitted a - /// (maybe deferred) "bad call" diagnostic. We use this to avoid emitting the - /// same deferred diag twice. - llvm::DenseSet<unsigned> LocsWithCUDACallDiags; + /// FunctionDecls plus raw encodings of SourceLocations for which + /// CheckCUDACall has emitted a (maybe deferred) "bad call" diagnostic. We + /// use this to avoid emitting the same deferred diag twice. + llvm::DenseSet<std::pair<FunctionDecl *, unsigned>> LocsWithCUDACallDiags; /// The set of CUDA functions that we've discovered must be emitted by tracing /// the call graph. Functions that we can tell a priori must be emitted Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=284643&r1=284642&r2=284643&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Oct 19 16:03:38 2016 @@ -714,20 +714,22 @@ bool Sema::CheckCUDACall(SourceLocation } }(); + if (DiagKind == CUDADiagBuilder::K_Nop) + return true; + // Avoid emitting this error twice for the same location. Using a hashtable // like this is unfortunate, but because we must continue parsing as normal // after encountering a deferred error, it's otherwise very tricky for us to // ensure that we only emit this deferred error once. - if (!LocsWithCUDACallDiags.insert(Loc.getRawEncoding()).second) + if (!LocsWithCUDACallDiags.insert({Caller, Loc.getRawEncoding()}).second) return true; - bool IsImmediateErr = - CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) + CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller); CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, Caller, *this) << Callee; - return !IsImmediateErr; + return DiagKind != CUDADiagBuilder::K_Immediate; } void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) { Added: cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu?rev=284643&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu (added) +++ cfe/trunk/test/SemaCUDA/bad-calls-on-same-line.cu Wed Oct 19 16:03:38 2016 @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +// The hd function template is instantiated three times. +// +// Two of those instantiations call a device function, which is an error when +// compiling for host. Clang should report both errors. + +#include "Inputs/cuda.h" + +template <typename T> +struct Selector {}; + +template <> +struct Selector<int> { + __host__ void f() {} +}; + +template <> +struct Selector<float> { + __device__ void f() {} // expected-note {{declared here}} +}; + +template <> +struct Selector<double> { + __device__ void f() {} // expected-note {{declared here}} +}; + +template <typename T> +inline __host__ __device__ void hd() { + Selector<T>().f(); + // expected-error@-1 {{reference to __device__ function}} + // expected-error@-2 {{reference to __device__ function}} +} + +void host_fn() { + hd<int>(); + hd<double>(); // expected-note {{function template specialization 'hd<double>'}} + hd<float>(); // expected-note {{function template specialization 'hd<float>'}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
