Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -10111,7 +10111,7 @@ /// compilation, this is currently only enabled for CUDA compilations. llvm::DenseMap, std::vector> - CUDADeferredDiags; + DeviceDeferredDiags; /// A pair of a canonical FunctionDecl and a SourceLocation. When used as the /// key in a hashtable, both the FD and location are hashed. @@ -10132,21 +10132,22 @@ /// map. llvm::DenseMap, /* Caller = */ FunctionDeclAndLoc> - CUDAKnownEmittedFns; + DeviceKnownEmittedFns; - /// A partial call graph maintained during CUDA compilation to support - /// deferred diagnostics. + /// A partial call graph maintained during CUDA/OpenMP device code compilation + /// to support deferred diagnostics. /// /// Functions are only added here if, at the time they're considered, they are /// not known-emitted. As soon as we discover that a function is /// known-emitted, we remove it and everything it transitively calls from this - /// set and add those functions to CUDAKnownEmittedFns. + /// set and add those functions to DeviceKnownEmittedFns. llvm::DenseMap, /* Callees = */ llvm::MapVector, SourceLocation>> - CUDACallGraph; + DeviceCallGraph; - /// Diagnostic builder for CUDA errors which may or may not be deferred. + /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be + /// deferred. /// /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch) /// which are not allowed to appear inside __device__ functions and are @@ -10160,7 +10161,7 @@ /// diagnostic, or no diagnostic at all, according to an argument you pass to /// its constructor, thus simplifying the process of creating these "maybe /// deferred" diagnostics. - class CUDADiagBuilder { + class DeviceDiagBuilder { public: enum Kind { /// Emit no diagnostics. @@ -10177,25 +10178,25 @@ K_Deferred }; - CUDADiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID, - FunctionDecl *Fn, Sema &S); - ~CUDADiagBuilder(); + DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID, + FunctionDecl *Fn, Sema &S); + ~DeviceDiagBuilder(); /// Convertible to bool: True if we immediately emitted an error, false if /// we didn't emit an error or we created a deferred error. /// /// Example usage: /// - /// if (CUDADiagBuilder(...) << foo << bar) + /// if (DeviceDiagBuilder(...) << foo << bar) /// return ExprError(); /// /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably - /// want to use these instead of creating a CUDADiagBuilder yourself. + /// want to use these instead of creating a DeviceDiagBuilder yourself. operator bool() const { return ImmediateDiag.hasValue(); } template - friend const CUDADiagBuilder &operator<<(const CUDADiagBuilder &Diag, - const T &Value) { + friend const DeviceDiagBuilder &operator<<(const DeviceDiagBuilder &Diag, + const T &Value) { if (Diag.ImmediateDiag.hasValue()) *Diag.ImmediateDiag << Value; else if (Diag.PartialDiag.hasValue()) @@ -10216,7 +10217,15 @@ llvm::Optional PartialDiag; }; - /// Creates a CUDADiagBuilder that emits the diagnostic if the current context + /// Indicate that this function (and thus everything it transtively calls) + /// will be codegen'ed, and emit any deferred diagnostics on this function and + /// its (transitive) callees. + void markKnownEmitted( + Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee, + SourceLocation OrigLoc, + const llvm::function_ref IsKnownEmitted); + + /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context /// is "used as device code". /// /// - If CurContext is a __host__ function, does not emit any diagnostics. @@ -10232,13 +10241,13 @@ /// if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget()) /// return ExprError(); /// // Otherwise, continue parsing as normal. - CUDADiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); + DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); - /// Creates a CUDADiagBuilder that emits the diagnostic if the current context + /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context /// is "used as host code". /// /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched. - CUDADiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID); + DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID); enum CUDAFunctionTarget { CFT_Device, Index: lib/Sema/Sema.cpp =================================================================== --- lib/Sema/Sema.cpp +++ lib/Sema/Sema.cpp @@ -1325,6 +1325,168 @@ return Builder; } +// Print notes showing how we can reach FD starting from an a priori +// known-callable function. +static void emitCallStackNotes(Sema &S, FunctionDecl *FD) { + auto FnIt = S.DeviceKnownEmittedFns.find(FD); + while (FnIt != S.DeviceKnownEmittedFns.end()) { + DiagnosticBuilder Builder( + S.Diags.Report(FnIt->second.Loc, diag::note_called_by)); + Builder << FnIt->second.FD; + Builder.setForceEmit(); + + FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD); + } +} + +// Emit any deferred diagnostics for FD and erase them from the map in which +// they're stored. +static void emitDeferredDiags(Sema &S, FunctionDecl *FD) { + auto It = S.DeviceDeferredDiags.find(FD); + if (It == S.DeviceDeferredDiags.end()) + return; + bool HasWarningOrError = false; + for (PartialDiagnosticAt &PDAt : It->second) { + const SourceLocation &Loc = PDAt.first; + const PartialDiagnostic &PD = PDAt.second; + HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel( + PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning; + DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID())); + Builder.setForceEmit(); + PD.Emit(Builder); + } + S.DeviceDeferredDiags.erase(It); + + // FIXME: Should this be called after every warning/error emitted in the loop + // above, instead of just once per function? That would be consistent with + // how we handle immediate errors, but it also seems like a bit much. + if (HasWarningOrError) + emitCallStackNotes(S, FD); +} + +// In CUDA, there are some constructs which may appear in semantically-valid +// code, but trigger errors if we ever generate code for the function in which +// they appear. Essentially every construct you're not allowed to use on the +// device falls into this category, because you are allowed to use these +// constructs in a __host__ __device__ function, but only if that function is +// never codegen'ed on the device. +// +// To handle semantic checking for these constructs, we keep track of the set of +// functions we know will be emitted, either because we could tell a priori that +// they would be emitted, or because they were transitively called by a +// known-emitted function. +// +// We also keep a partial call graph of which not-known-emitted functions call +// which other not-known-emitted functions. +// +// When we see something which is illegal if the current function is emitted +// (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or +// CheckCUDACall), we first check if the current function is known-emitted. If +// so, we immediately output the diagnostic. +// +// Otherwise, we "defer" the diagnostic. It sits in Sema::DeviceDeferredDiags +// until we discover that the function is known-emitted, at which point we take +// it out of this map and emit the diagnostic. + +Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc, + unsigned DiagID, FunctionDecl *Fn, + Sema &S) + : S(S), Loc(Loc), DiagID(DiagID), Fn(Fn), + ShowCallStack(K == K_ImmediateWithCallStack || K == K_Deferred) { + switch (K) { + case K_Nop: + break; + case K_Immediate: + case K_ImmediateWithCallStack: + ImmediateDiag.emplace(S.Diag(Loc, DiagID)); + break; + case K_Deferred: + assert(Fn && "Must have a function to attach the deferred diag to."); + PartialDiag.emplace(S.PDiag(DiagID)); + break; + } +} + +Sema::DeviceDiagBuilder::~DeviceDiagBuilder() { + if (ImmediateDiag) { + // Emit our diagnostic and, if it was a warning or error, output a callstack + // if Fn isn't a priori known-emitted. + bool IsWarningOrError = S.getDiagnostics().getDiagnosticLevel( + DiagID, Loc) >= DiagnosticsEngine::Warning; + ImmediateDiag.reset(); // Emit the immediate diag. + if (IsWarningOrError && ShowCallStack) + emitCallStackNotes(S, Fn); + } else if (PartialDiag) { + assert(ShowCallStack && "Must always show call stack for deferred diags."); + S.DeviceDeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)}); + } +} + +// Indicate that this function (and thus everything it transtively calls) will +// be codegen'ed, and emit any deferred diagnostics on this function and its +// (transitive) callees. +void Sema::markKnownEmitted( + Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee, + SourceLocation OrigLoc, + const llvm::function_ref IsKnownEmitted) { + // Nothing to do if we already know that FD is emitted. + if (IsKnownEmitted(S, OrigCallee)) { + assert(!S.DeviceCallGraph.count(OrigCallee)); + return; + } + + // We've just discovered that OrigCallee is known-emitted. Walk our call + // graph to see what else we can now discover also must be emitted. + + struct CallInfo { + FunctionDecl *Caller; + FunctionDecl *Callee; + SourceLocation Loc; + }; + llvm::SmallVector Worklist = {{OrigCaller, OrigCallee, OrigLoc}}; + llvm::SmallSet, 4> Seen; + Seen.insert(OrigCallee); + while (!Worklist.empty()) { + CallInfo C = Worklist.pop_back_val(); + assert(!IsKnownEmitted(S, C.Callee) && + "Worklist should not contain known-emitted functions."); + S.DeviceKnownEmittedFns[C.Callee] = {C.Caller, C.Loc}; + emitDeferredDiags(S, C.Callee); + + // If this is a template instantiation, explore its callgraph as well: + // Non-dependent calls are part of the template's callgraph, while dependent + // calls are part of to the instantiation's call graph. + if (auto *Templ = C.Callee->getPrimaryTemplate()) { + FunctionDecl *TemplFD = Templ->getAsFunction(); + if (!Seen.count(TemplFD) && !S.DeviceKnownEmittedFns.count(TemplFD)) { + Seen.insert(TemplFD); + Worklist.push_back( + {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc}); + } + } + + // Add all functions called by Callee to our worklist. + auto CGIt = S.DeviceCallGraph.find(C.Callee); + if (CGIt == S.DeviceCallGraph.end()) + continue; + + for (std::pair, SourceLocation> FDLoc : + CGIt->second) { + FunctionDecl *NewCallee = FDLoc.first; + SourceLocation CallLoc = FDLoc.second; + if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee)) + continue; + Seen.insert(NewCallee); + Worklist.push_back( + {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc}); + } + + // C.Callee is now known-emitted, so we no longer need to maintain its list + // of callees in DeviceCallGraph. + S.DeviceCallGraph.erase(CGIt); + } +} + /// Looks through the macro-expansion chain for the given /// location, looking for a macro expansion with the given name. /// If one is found, returns true and sets the location to that Index: lib/Sema/SemaCUDA.cpp =================================================================== --- lib/Sema/SemaCUDA.cpp +++ lib/Sema/SemaCUDA.cpp @@ -586,78 +586,6 @@ NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); } -// In CUDA, there are some constructs which may appear in semantically-valid -// code, but trigger errors if we ever generate code for the function in which -// they appear. Essentially every construct you're not allowed to use on the -// device falls into this category, because you are allowed to use these -// constructs in a __host__ __device__ function, but only if that function is -// never codegen'ed on the device. -// -// To handle semantic checking for these constructs, we keep track of the set of -// functions we know will be emitted, either because we could tell a priori that -// they would be emitted, or because they were transitively called by a -// known-emitted function. -// -// We also keep a partial call graph of which not-known-emitted functions call -// which other not-known-emitted functions. -// -// When we see something which is illegal if the current function is emitted -// (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or -// CheckCUDACall), we first check if the current function is known-emitted. If -// so, we immediately output the diagnostic. -// -// Otherwise, we "defer" the diagnostic. It sits in Sema::CUDADeferredDiags -// until we discover that the function is known-emitted, at which point we take -// it out of this map and emit the diagnostic. - -Sema::CUDADiagBuilder::CUDADiagBuilder(Kind K, SourceLocation Loc, - unsigned DiagID, FunctionDecl *Fn, - Sema &S) - : S(S), Loc(Loc), DiagID(DiagID), Fn(Fn), - ShowCallStack(K == K_ImmediateWithCallStack || K == K_Deferred) { - switch (K) { - case K_Nop: - break; - case K_Immediate: - case K_ImmediateWithCallStack: - ImmediateDiag.emplace(S.Diag(Loc, DiagID)); - break; - case K_Deferred: - assert(Fn && "Must have a function to attach the deferred diag to."); - PartialDiag.emplace(S.PDiag(DiagID)); - break; - } -} - -// Print notes showing how we can reach FD starting from an a priori -// known-callable function. -static void EmitCallStackNotes(Sema &S, FunctionDecl *FD) { - auto FnIt = S.CUDAKnownEmittedFns.find(FD); - while (FnIt != S.CUDAKnownEmittedFns.end()) { - DiagnosticBuilder Builder( - S.Diags.Report(FnIt->second.Loc, diag::note_called_by)); - Builder << FnIt->second.FD; - Builder.setForceEmit(); - - FnIt = S.CUDAKnownEmittedFns.find(FnIt->second.FD); - } -} - -Sema::CUDADiagBuilder::~CUDADiagBuilder() { - if (ImmediateDiag) { - // Emit our diagnostic and, if it was a warning or error, output a callstack - // if Fn isn't a priori known-emitted. - bool IsWarningOrError = S.getDiagnostics().getDiagnosticLevel( - DiagID, Loc) >= DiagnosticsEngine::Warning; - ImmediateDiag.reset(); // Emit the immediate diag. - if (IsWarningOrError && ShowCallStack) - EmitCallStackNotes(S, Fn); - } else if (PartialDiag) { - assert(ShowCallStack && "Must always show call stack for deferred diags."); - S.CUDADeferredDiags[Fn].push_back({Loc, std::move(*PartialDiag)}); - } -} - // Do we know that we will eventually codegen the given function? static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) { // Templates are emitted when they're instantiated. @@ -689,147 +617,59 @@ // Otherwise, the function is known-emitted if it's in our set of // known-emitted functions. - return S.CUDAKnownEmittedFns.count(FD) > 0; + return S.DeviceKnownEmittedFns.count(FD) > 0; } -Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, - unsigned DiagID) { +Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - CUDADiagBuilder::Kind DiagKind = [&] { + DeviceDiagBuilder::Kind DiagKind = [this] { switch (CurrentCUDATarget()) { case CFT_Global: case CFT_Device: - return CUDADiagBuilder::K_Immediate; + return DeviceDiagBuilder::K_Immediate; case CFT_HostDevice: // An HD function counts as host code if we're compiling for host, and // device code if we're compiling for device. Defer any errors in device // mode until the function is known-emitted. if (getLangOpts().CUDAIsDevice) { return IsKnownEmitted(*this, dyn_cast(CurContext)) - ? CUDADiagBuilder::K_ImmediateWithCallStack - : CUDADiagBuilder::K_Deferred; + ? DeviceDiagBuilder::K_ImmediateWithCallStack + : DeviceDiagBuilder::K_Deferred; } - return CUDADiagBuilder::K_Nop; + return DeviceDiagBuilder::K_Nop; default: - return CUDADiagBuilder::K_Nop; + return DeviceDiagBuilder::K_Nop; } }(); - return CUDADiagBuilder(DiagKind, Loc, DiagID, - dyn_cast(CurContext), *this); + return DeviceDiagBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); } -Sema::CUDADiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc, - unsigned DiagID) { +Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc, + unsigned DiagID) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - CUDADiagBuilder::Kind DiagKind = [&] { + DeviceDiagBuilder::Kind DiagKind = [this] { switch (CurrentCUDATarget()) { case CFT_Host: - return CUDADiagBuilder::K_Immediate; + return DeviceDiagBuilder::K_Immediate; case CFT_HostDevice: // An HD function counts as host code if we're compiling for host, and // device code if we're compiling for device. Defer any errors in device // mode until the function is known-emitted. if (getLangOpts().CUDAIsDevice) - return CUDADiagBuilder::K_Nop; + return DeviceDiagBuilder::K_Nop; return IsKnownEmitted(*this, dyn_cast(CurContext)) - ? CUDADiagBuilder::K_ImmediateWithCallStack - : CUDADiagBuilder::K_Deferred; + ? DeviceDiagBuilder::K_ImmediateWithCallStack + : DeviceDiagBuilder::K_Deferred; default: - return CUDADiagBuilder::K_Nop; + return DeviceDiagBuilder::K_Nop; } }(); - return CUDADiagBuilder(DiagKind, Loc, DiagID, - dyn_cast(CurContext), *this); -} - -// Emit any deferred diagnostics for FD and erase them from the map in which -// they're stored. -static void EmitDeferredDiags(Sema &S, FunctionDecl *FD) { - auto It = S.CUDADeferredDiags.find(FD); - if (It == S.CUDADeferredDiags.end()) - return; - bool HasWarningOrError = false; - for (PartialDiagnosticAt &PDAt : It->second) { - const SourceLocation &Loc = PDAt.first; - const PartialDiagnostic &PD = PDAt.second; - HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel( - PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning; - DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID())); - Builder.setForceEmit(); - PD.Emit(Builder); - } - S.CUDADeferredDiags.erase(It); - - // FIXME: Should this be called after every warning/error emitted in the loop - // above, instead of just once per function? That would be consistent with - // how we handle immediate errors, but it also seems like a bit much. - if (HasWarningOrError) - EmitCallStackNotes(S, FD); -} - -// Indicate that this function (and thus everything it transtively calls) will -// be codegen'ed, and emit any deferred diagnostics on this function and its -// (transitive) callees. -static void MarkKnownEmitted(Sema &S, FunctionDecl *OrigCaller, - FunctionDecl *OrigCallee, SourceLocation OrigLoc) { - // Nothing to do if we already know that FD is emitted. - if (IsKnownEmitted(S, OrigCallee)) { - assert(!S.CUDACallGraph.count(OrigCallee)); - return; - } - - // We've just discovered that OrigCallee is known-emitted. Walk our call - // graph to see what else we can now discover also must be emitted. - - struct CallInfo { - FunctionDecl *Caller; - FunctionDecl *Callee; - SourceLocation Loc; - }; - llvm::SmallVector Worklist = {{OrigCaller, OrigCallee, OrigLoc}}; - llvm::SmallSet, 4> Seen; - Seen.insert(OrigCallee); - while (!Worklist.empty()) { - CallInfo C = Worklist.pop_back_val(); - assert(!IsKnownEmitted(S, C.Callee) && - "Worklist should not contain known-emitted functions."); - S.CUDAKnownEmittedFns[C.Callee] = {C.Caller, C.Loc}; - EmitDeferredDiags(S, C.Callee); - - // If this is a template instantiation, explore its callgraph as well: - // Non-dependent calls are part of the template's callgraph, while dependent - // calls are part of to the instantiation's call graph. - if (auto *Templ = C.Callee->getPrimaryTemplate()) { - FunctionDecl *TemplFD = Templ->getAsFunction(); - if (!Seen.count(TemplFD) && !S.CUDAKnownEmittedFns.count(TemplFD)) { - Seen.insert(TemplFD); - Worklist.push_back( - {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc}); - } - } - - // Add all functions called by Callee to our worklist. - auto CGIt = S.CUDACallGraph.find(C.Callee); - if (CGIt == S.CUDACallGraph.end()) - continue; - - for (std::pair, SourceLocation> FDLoc : - CGIt->second) { - FunctionDecl *NewCallee = FDLoc.first; - SourceLocation CallLoc = FDLoc.second; - if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee)) - continue; - Seen.insert(NewCallee); - Worklist.push_back( - {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc}); - } - - // C.Callee is now known-emitted, so we no longer need to maintain its list - // of callees in CUDACallGraph. - S.CUDACallGraph.erase(CGIt); - } + return DeviceDiagBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); } bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { @@ -848,7 +688,7 @@ // Host-side references to a __global__ function refer to the stub, so the // function itself is never emitted and therefore should not be marked. if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global) - MarkKnownEmitted(*this, Caller, Callee, Loc); + markKnownEmitted(*this, Caller, Callee, Loc, IsKnownEmitted); } else { // If we have // host fn calls kernel fn calls host+device, @@ -857,26 +697,27 @@ // that, when compiling for host, only HD functions actually called from the // host get marked as known-emitted. if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global) - CUDACallGraph[Caller].insert({Callee, Loc}); + DeviceCallGraph[Caller].insert({Callee, Loc}); } - CUDADiagBuilder::Kind DiagKind = [&] { + DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee, + CallerKnownEmitted] { switch (IdentifyCUDAPreference(Caller, Callee)) { case CFP_Never: - return CUDADiagBuilder::K_Immediate; + return DeviceDiagBuilder::K_Immediate; case CFP_WrongSide: assert(Caller && "WrongSide calls require a non-null caller"); // If we know the caller will be emitted, we know this wrong-side call // will be emitted, so it's an immediate error. Otherwise, defer the // error until we know the caller is emitted. - return CallerKnownEmitted ? CUDADiagBuilder::K_ImmediateWithCallStack - : CUDADiagBuilder::K_Deferred; + return CallerKnownEmitted ? DeviceDiagBuilder::K_ImmediateWithCallStack + : DeviceDiagBuilder::K_Deferred; default: - return CUDADiagBuilder::K_Nop; + return DeviceDiagBuilder::K_Nop; } }(); - if (DiagKind == CUDADiagBuilder::K_Nop) + if (DiagKind == DeviceDiagBuilder::K_Nop) return true; // Avoid emitting this error twice for the same location. Using a hashtable @@ -886,13 +727,13 @@ if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second) return true; - CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) + DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller); - CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, - Caller, *this) + DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, + Caller, *this) << Callee; - return DiagKind != CUDADiagBuilder::K_Immediate && - DiagKind != CUDADiagBuilder::K_ImmediateWithCallStack; + return DiagKind != DeviceDiagBuilder::K_Immediate && + DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack; } void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {