r284643 - [CUDA] Emit errors for wrong-side calls made on the same line as non-wrong-side calls.

[Justin Lebar via cfe-commits]

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 at -1 {{reference to __device__ function}}
+  // expected-error at -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>'}}
+}