maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA] Fix "declared here" note on deferred wrong-side errors. 

Previously we weren't deferring these "declared here" notes, which is
obviously wrong.

llvm-svn: 278767
This commit is contained in:
Justin Lebar 2016-08-16 00:48:21 +00:00
parent acc50c4334
commit 26bb31123a
3 changed files with 29 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
clang/lib/Sema/SemaCUDA.cpp
View File

@ -499,11 +499,16 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
if (Pref == Sema::CFP_WrongSide) {
// We have to do this odd dance to create our PartialDiagnostic because we
// want its storage to be allocated with operator new, not in an arena.
PartialDiagnostic PD{PartialDiagnostic::NullDiagnostic()};
PD.Reset(diag::err_ref_bad_target);
PD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
Caller->addDeferredDiag({Loc, std::move(PD)});
Diag(Callee->getLocation(), diag::note_previous_decl) << Callee;
PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()};
ErrPD.Reset(diag::err_ref_bad_target);
ErrPD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
Caller->addDeferredDiag({Loc, std::move(ErrPD)});
PartialDiagnostic NotePD{PartialDiagnostic::NullDiagnostic()};
NotePD.Reset(diag::note_previous_decl);
NotePD << Callee;
Caller->addDeferredDiag({Callee->getLocation(), std::move(NotePD)});
// This is not immediately an error, so return true. The deferred errors
// will be emitted if and when Caller is codegen'ed.
return true;

9
clang/test/SemaCUDA/call-device-fn-from-host.cu
View File

@ -6,10 +6,18 @@
#include "Inputs/cuda.h"
__device__ void device_fn() {}
// expected-note@-1 {{'device_fn' declared here}}
// expected-note@-2 {{'device_fn' declared here}}
// expected-note@-3 {{'device_fn' declared here}}
// expected-note@-4 {{'device_fn' declared here}}
// expected-note@-5 {{'device_fn' declared here}}
struct S {
__device__ S() {}
// expected-note@-1 {{'S' declared here}}
// expected-note@-2 {{'S' declared here}}
__device__ ~S() { device_fn(); }
// expected-note@-1 {{'~S' declared here}}
int x;
};
@ -24,6 +32,7 @@ struct T {
__host__ __device__ void hd3();
__device__ void d() {}
// expected-note@-1 {{'d' declared here}}
};
__host__ __device__ void T::hd3() {

10
clang/test/SemaCUDA/call-host-fn-from-device.cu
View File

@ -6,10 +6,19 @@
#include "Inputs/cuda.h"
extern "C" void host_fn() {}
// expected-note@-1 {{'host_fn' declared here}}
// expected-note@-2 {{'host_fn' declared here}}
// expected-note@-3 {{'host_fn' declared here}}
// expected-note@-4 {{'host_fn' declared here}}
// expected-note@-5 {{'host_fn' declared here}}
// expected-note@-6 {{'host_fn' declared here}}
struct S {
S() {}
// expected-note@-1 {{'S' declared here}}
// expected-note@-2 {{'S' declared here}}
~S() { host_fn(); }
// expected-note@-1 {{'~S' declared here}}
int x;
};
@ -24,6 +33,7 @@ struct T {
__host__ __device__ void hd3();
void h() {}
// expected-note@-1 {{'h' declared here}}
};
__host__ __device__ void T::hd3() {