/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp

Bug Summary

File:	clang/lib/CodeGen/CGCUDANV.cpp
Warning:	line 1032, column 13 Called C++ object pointer is null

Annotated Source Code

Press '?' to see keyboard shortcuts

Show analyzer invocation

clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -analyze -disable-free -disable-llvm-verifier -discard-value-names -main-file-name CGCUDANV.cpp -analyzer-store=region -analyzer-opt-analyze-nested-blocks -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -fhalf-no-semantic-interposition -mframe-pointer=none -relaxed-aliasing -fmath-errno -fno-rounding-math -mconstructor-aliases -munwind-tables -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -ffunction-sections -fdata-sections -fcoverage-compilation-dir=/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/build-llvm/tools/clang/lib/CodeGen -resource-dir /usr/lib/llvm-13/lib/clang/13.0.0 -D CLANG_ROUND_TRIP_CC1_ARGS=ON -D _DEBUG -D _GNU_SOURCE -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/build-llvm/tools/clang/lib/CodeGen -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/build-llvm/tools/clang/include -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/build-llvm/include -I /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/llvm/include -U NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/c++/6.3.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/x86_64-linux-gnu/c++/6.3.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../include/c++/6.3.0/backward -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/6.3.0/../../../../x86_64-linux-gnu/include -internal-isystem /usr/lib/llvm-13/lib/clang/13.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -O2 -Wno-unused-parameter -Wwrite-strings -Wno-missing-field-initializers -Wno-long-long -Wno-maybe-uninitialized -Wno-comment -std=c++14 -fdeprecated-macro -fdebug-compilation-dir=/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/build-llvm/tools/clang/lib/CodeGen -fdebug-prefix-map=/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4=. -ferror-limit 19 -fvisibility-inlines-hidden -stack-protector 2 -fgnuc-version=4.2.1 -vectorize-loops -vectorize-slp -analyzer-output=html -analyzer-config stable-report-filename=true -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/scan-build-2021-04-05-202135-9119-1 -x c++ /build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp

/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp

→

1//===----- CGCUDANV.cpp - Interface to NVIDIA CUDA Runtime ----------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This provides a class for CUDA code generation targeting the NVIDIA CUDA
10// runtime library.
11//
12//===----------------------------------------------------------------------===//

14#include "CGCUDARuntime.h"
15#include "CGCXXABI.h"
16#include "CodeGenFunction.h"
17#include "CodeGenModule.h"
18#include "clang/AST/Decl.h"
19#include "clang/Basic/Cuda.h"
20#include "clang/CodeGen/CodeGenABITypes.h"
21#include "clang/CodeGen/ConstantInitBuilder.h"
22#include "llvm/IR/BasicBlock.h"
23#include "llvm/IR/Constants.h"
24#include "llvm/IR/DerivedTypes.h"
25#include "llvm/IR/ReplaceConstant.h"
26#include "llvm/Support/Format.h"

28using namespace clang;
29using namespace CodeGen;

31namespace {
32constexpr unsigned CudaFatMagic = 0x466243b1;
33constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF"

35class CGNVCUDARuntime : public CGCUDARuntime {

37private:
llvm::IntegerType *IntTy, *SizeTy;
llvm::Type *VoidTy;
llvm::PointerType *CharPtrTy, *VoidPtrTy, *VoidPtrPtrTy;

/// Convenience reference to LLVM Context
llvm::LLVMContext &Context;
/// Convenience reference to the current module
llvm::Module &TheModule;
/// Keeps track of kernel launch stubs and handles emitted in this module
struct KernelInfo {
  llvm::Function *Kernel; // stub function to help launch kernel
  const Decl *D;
};
llvm::SmallVector<KernelInfo, 16> EmittedKernels;
// Map a device stub function to a symbol for identifying kernel in host code.
// For CUDA, the symbol for identifying the kernel is the same as the device
// stub function. For HIP, they are different.
llvm::DenseMap<llvm::Function *, llvm::GlobalValue *> KernelHandles;
// Map a kernel handle to the kernel stub.
llvm::DenseMap<llvm::GlobalValue *, llvm::Function *> KernelStubs;
struct VarInfo {
  llvm::GlobalVariable *Var;
  const VarDecl *D;
  DeviceVarFlags Flags;
};
llvm::SmallVector<VarInfo, 16> DeviceVars;
/// Keeps track of variable containing handle of GPU binary. Populated by
/// ModuleCtorFunction() and used to create corresponding cleanup calls in
/// ModuleDtorFunction()
llvm::GlobalVariable *GpuBinaryHandle = nullptr;
/// Whether we generate relocatable device code.
bool RelocatableDeviceCode;
/// Mangle context for device.
std::unique_ptr<MangleContext> DeviceMC;

llvm::FunctionCallee getSetupArgumentFn() const;
llvm::FunctionCallee getLaunchFn() const;

llvm::FunctionType *getRegisterGlobalsFnTy() const;
llvm::FunctionType *getCallbackFnTy() const;
llvm::FunctionType *getRegisterLinkedBinaryFnTy() const;
std::string addPrefixToName(StringRef FuncName) const;
std::string addUnderscoredPrefixToName(StringRef FuncName) const;

/// Creates a function to register all kernel stubs generated in this module.
llvm::Function *makeRegisterGlobalsFn();

/// Helper function that generates a constant string and returns a pointer to
/// the start of the string.  The result of this function can be used anywhere
/// where the C code specifies const char*.
llvm::Constant *makeConstantString(const std::string &Str,
                                   const std::string &Name = "",
                                   const std::string &SectionName = "",
                                   unsigned Alignment = 0) {
  llvm::Constant *Zeros[] = {llvm::ConstantInt::get(SizeTy, 0),
                             llvm::ConstantInt::get(SizeTy, 0)};
  auto ConstStr = CGM.GetAddrOfConstantCString(Str, Name.c_str());
  llvm::GlobalVariable *GV =
      cast<llvm::GlobalVariable>(ConstStr.getPointer());
  if (!SectionName.empty()) {
    GV->setSection(SectionName);
    // Mark the address as used which make sure that this section isn't
    // merged and we will really have it in the object file.
    GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::None);
  }
  if (Alignment)
    GV->setAlignment(llvm::Align(Alignment));

  return llvm::ConstantExpr::getGetElementPtr(ConstStr.getElementType(),
                                              ConstStr.getPointer(), Zeros);
}

/// Helper function that generates an empty dummy function returning void.
llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) {
  assert(FnTy->getReturnType()->isVoidTy() &&((FnTy->getReturnType()->isVoidTy() && "Can only generate dummy functions returning void!"
) ? static_cast<void> (0) : __assert_fail ("FnTy->getReturnType()->isVoidTy() && \"Can only generate dummy functions returning void!\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 113, __PRETTY_FUNCTION__))
         "Can only generate dummy functions returning void!")((FnTy->getReturnType()->isVoidTy() && "Can only generate dummy functions returning void!"
) ? static_cast<void> (0) : __assert_fail ("FnTy->getReturnType()->isVoidTy() && \"Can only generate dummy functions returning void!\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 113, __PRETTY_FUNCTION__));
  llvm::Function *DummyFunc = llvm::Function::Create(
      FnTy, llvm::GlobalValue::InternalLinkage, "dummy", &TheModule);

  llvm::BasicBlock *DummyBlock =
      llvm::BasicBlock::Create(Context, "", DummyFunc);
  CGBuilderTy FuncBuilder(CGM, Context);
  FuncBuilder.SetInsertPoint(DummyBlock);
  FuncBuilder.CreateRetVoid();

  return DummyFunc;
}

void emitDeviceStubBodyLegacy(CodeGenFunction &CGF, FunctionArgList &Args);
void emitDeviceStubBodyNew(CodeGenFunction &CGF, FunctionArgList &Args);
std::string getDeviceSideName(const NamedDecl *ND) override;

void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
                       bool Extern, bool Constant) {
  DeviceVars.push_back({&Var,
                        VD,
                        {DeviceVarFlags::Variable, Extern, Constant,
                         VD->hasAttr<HIPManagedAttr>(),
                         /*Normalized*/ false, 0}});
}
void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
                        bool Extern, int Type) {
  DeviceVars.push_back({&Var,
                        VD,
                        {DeviceVarFlags::Surface, Extern, /*Constant*/ false,
                         /*Managed*/ false,
                         /*Normalized*/ false, Type}});
}
void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
                       bool Extern, int Type, bool Normalized) {
  DeviceVars.push_back({&Var,
                        VD,
                        {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
                         /*Managed*/ false, Normalized, Type}});
}

/// Creates module constructor function
llvm::Function *makeModuleCtorFunction();
/// Creates module destructor function
llvm::Function *makeModuleDtorFunction();
/// Transform managed variables for device compilation.
void transformManagedVars();

161public:
CGNVCUDARuntime(CodeGenModule &CGM);

llvm::GlobalValue *getKernelHandle(llvm::Function *F, GlobalDecl GD) override;
llvm::Function *getKernelStub(llvm::GlobalValue *Handle) override {
  auto Loc = KernelStubs.find(Handle);
  assert(Loc != KernelStubs.end())((Loc != KernelStubs.end()) ? static_cast<void> (0) : __assert_fail
 ("Loc != KernelStubs.end()", "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 167, __PRETTY_FUNCTION__));
  return Loc->second;
}
void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
void handleVarRegistration(const VarDecl *VD,
                           llvm::GlobalVariable &Var) override;
void
internalizeDeviceSideVar(const VarDecl *D,
                         llvm::GlobalValue::LinkageTypes &Linkage) override;

llvm::Function *finalizeModule() override;
178};

180}

182std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
if (CGM.getLangOpts().HIP)
  return ((Twine("hip") + Twine(FuncName)).str());
return ((Twine("cuda") + Twine(FuncName)).str());
186}
187std::string
188CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
if (CGM.getLangOpts().HIP)
  return ((Twine("__hip") + Twine(FuncName)).str());
return ((Twine("__cuda") + Twine(FuncName)).str());
192}

194CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
  : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
    TheModule(CGM.getModule()),
    RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode),
    DeviceMC(CGM.getContext().createMangleContext(
        CGM.getContext().getAuxTargetInfo())) {
CodeGen::CodeGenTypes &Types = CGM.getTypes();
ASTContext &Ctx = CGM.getContext();

IntTy = CGM.IntTy;
SizeTy = CGM.SizeTy;
VoidTy = CGM.VoidTy;

CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
VoidPtrPtrTy = VoidPtrTy->getPointerTo();
if (CGM.getContext().getAuxTargetInfo()) {
  // If the host and device have different C++ ABIs, mark it as the device
  // mangle context so that the mangling needs to retrieve the additonal
  // device lambda mangling number instead of the regular host one.
  DeviceMC->setDeviceMangleContext(
      CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
      CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily());
}
218}

220llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
// cudaError_t cudaSetupArgument(void *, size_t, size_t)
llvm::Type *Params[] = {VoidPtrTy, SizeTy, SizeTy};
return CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(IntTy, Params, false),
    addPrefixToName("SetupArgument"));
226}

228llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const {
if (CGM.getLangOpts().HIP) {
  // hipError_t hipLaunchByPtr(char *);
  return CGM.CreateRuntimeFunction(
      llvm::FunctionType::get(IntTy, CharPtrTy, false), "hipLaunchByPtr");
} else {
  // cudaError_t cudaLaunch(char *);
  return CGM.CreateRuntimeFunction(
      llvm::FunctionType::get(IntTy, CharPtrTy, false), "cudaLaunch");
}
238}

240llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
return llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false);
242}

244llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
return llvm::FunctionType::get(VoidTy, VoidPtrTy, false);
246}

248llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
auto CallbackFnTy = getCallbackFnTy();
auto RegisterGlobalsFnTy = getRegisterGlobalsFnTy();
llvm::Type *Params[] = {RegisterGlobalsFnTy->getPointerTo(), VoidPtrTy,
                        VoidPtrTy, CallbackFnTy->getPointerTo()};
return llvm::FunctionType::get(VoidTy, Params, false);
254}

256std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
GlobalDecl GD;
// D could be either a kernel or a variable.
if (auto *FD = dyn_cast<FunctionDecl>(ND))
  GD = GlobalDecl(FD, KernelReferenceKind::Kernel);
else
  GD = GlobalDecl(ND);
std::string DeviceSideName;
MangleContext *MC;
if (CGM.getLangOpts().CUDAIsDevice)
  MC = &CGM.getCXXABI().getMangleContext();
else
  MC = DeviceMC.get();
if (MC->shouldMangleDeclName(ND)) {
  SmallString<256> Buffer;
  llvm::raw_svector_ostream Out(Buffer);
  MC->mangleName(GD, Out);
  DeviceSideName = std::string(Out.str());
} else
  DeviceSideName = std::string(ND->getIdentifier()->getName());

// Make unique name for device side static file-scope variable for HIP.
if (CGM.getContext().shouldExternalizeStaticVar(ND) &&
    CGM.getLangOpts().GPURelocatableDeviceCode &&
    !CGM.getLangOpts().CUID.empty()) {
  SmallString<256> Buffer;
  llvm::raw_svector_ostream Out(Buffer);
  Out << DeviceSideName;
  CGM.printPostfixForExternalizedStaticVar(Out);
  DeviceSideName = std::string(Out.str());
}
return DeviceSideName;
288}

290void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
                                   FunctionArgList &Args) {
EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
if (auto *GV = dyn_cast<llvm::GlobalVariable>(KernelHandles[CGF.CurFn])) {
  GV->setLinkage(CGF.CurFn->getLinkage());
  GV->setInitializer(CGF.CurFn);
}
if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                       CudaFeature::CUDA_USES_NEW_LAUNCH) ||
    (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI))
  emitDeviceStubBodyNew(CGF, Args);
else
  emitDeviceStubBodyLegacy(CGF, Args);
303}

305// CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
306// array and kernels are launched using cudaLaunchKernel().
307void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
                                          FunctionArgList &Args) {
// Build the shadow stack entry at the very start of the function.

// Calculate amount of space we will need for all arguments.  If we have no
// args, allocate a single pointer so we still have a valid pointer to the
// argument array that we can pass to runtime, even if it will be unused.
Address KernelArgs = CGF.CreateTempAlloca(
    VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args",
    llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
// Store pointers to the arguments in a locally allocated launch_args.
for (unsigned i = 0; i < Args.size(); ++i) {
  llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
  llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, VoidPtrTy);
  CGF.Builder.CreateDefaultAlignedStore(
      VoidVarPtr, CGF.Builder.CreateConstGEP1_32(KernelArgs.getPointer(), i));
}

llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");

// Lookup cudaLaunchKernel/hipLaunchKernel function.
// cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
//                              void **args, size_t sharedMem,
//                              cudaStream_t stream);
// hipError_t hipLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
//                            void **args, size_t sharedMem,
//                            hipStream_t stream);
TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
auto LaunchKernelName = addPrefixToName("LaunchKernel");
IdentifierInfo &cudaLaunchKernelII =
    CGM.getContext().Idents.get(LaunchKernelName);
FunctionDecl *cudaLaunchKernelFD = nullptr;
for (const auto &Result : DC->lookup(&cudaLaunchKernelII)) {
  if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
    cudaLaunchKernelFD = FD;
}

if (cudaLaunchKernelFD == nullptr) {
  CGM.Error(CGF.CurFuncDecl->getLocation(),
            "Can't find declaration for " + LaunchKernelName);
  return;
}
// Create temporary dim3 grid_dim, block_dim.
ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1);
QualType Dim3Ty = GridDimParam->getType();
Address GridDim =
    CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim");
Address BlockDim =
    CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
Address ShmemSize =
    CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
Address Stream =
    CGF.CreateTempAlloca(VoidPtrTy, CGM.getPointerAlign(), "stream");
llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(IntTy,
                            {/*gridDim=*/GridDim.getType(),
                             /*blockDim=*/BlockDim.getType(),
                             /*ShmemSize=*/ShmemSize.getType(),
                             /*Stream=*/Stream.getType()},
                            /*isVarArg=*/false),
    addUnderscoredPrefixToName("PopCallConfiguration"));

CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn,
                            {GridDim.getPointer(), BlockDim.getPointer(),
                             ShmemSize.getPointer(), Stream.getPointer()});

// Emit the call to cudaLaunch
llvm::Value *Kernel =
    CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn], VoidPtrTy);
CallArgList LaunchKernelArgs;
LaunchKernelArgs.add(RValue::get(Kernel),
                     cudaLaunchKernelFD->getParamDecl(0)->getType());
LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty);
LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty);
LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()),
                     cudaLaunchKernelFD->getParamDecl(3)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)),
                     cudaLaunchKernelFD->getParamDecl(4)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)),
                     cudaLaunchKernelFD->getParamDecl(5)->getType());

QualType QT = cudaLaunchKernelFD->getType();
QualType CQT = QT.getCanonicalType();
llvm::Type *Ty = CGM.getTypes().ConvertType(CQT);
llvm::FunctionType *FTy = dyn_cast<llvm::FunctionType>(Ty);

const CGFunctionInfo &FI =
    CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD);
llvm::FunctionCallee cudaLaunchKernelFn =
    CGM.CreateRuntimeFunction(FTy, LaunchKernelName);
CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(),
             LaunchKernelArgs);
CGF.EmitBranch(EndBlock);

CGF.EmitBlock(EndBlock);
403}

405void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
                                             FunctionArgList &Args) {
// Emit a call to cudaSetupArgument for each arg in Args.
llvm::FunctionCallee cudaSetupArgFn = getSetupArgumentFn();
llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
CharUnits Offset = CharUnits::Zero();
for (const VarDecl *A : Args) {
  auto TInfo = CGM.getContext().getTypeInfoInChars(A->getType());
  Offset = Offset.alignTo(TInfo.Align);
  llvm::Value *Args[] = {
      CGF.Builder.CreatePointerCast(CGF.GetAddrOfLocalVar(A).getPointer(),
                                    VoidPtrTy),
      llvm::ConstantInt::get(SizeTy, TInfo.Width.getQuantity()),
      llvm::ConstantInt::get(SizeTy, Offset.getQuantity()),
  };
  llvm::CallBase *CB = CGF.EmitRuntimeCallOrInvoke(cudaSetupArgFn, Args);
  llvm::Constant *Zero = llvm::ConstantInt::get(IntTy, 0);
  llvm::Value *CBZero = CGF.Builder.CreateICmpEQ(CB, Zero);
  llvm::BasicBlock *NextBlock = CGF.createBasicBlock("setup.next");
  CGF.Builder.CreateCondBr(CBZero, NextBlock, EndBlock);
  CGF.EmitBlock(NextBlock);
  Offset += TInfo.Width;
}

// Emit the call to cudaLaunch
llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
llvm::Value *Arg =
    CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn], CharPtrTy);
CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg);
CGF.EmitBranch(EndBlock);

CGF.EmitBlock(EndBlock);
437}

439// Replace the original variable Var with the address loaded from variable
440// ManagedVar populated by HIP runtime.
441static void replaceManagedVar(llvm::GlobalVariable *Var,
                            llvm::GlobalVariable *ManagedVar) {
SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList;
for (auto &&VarUse : Var->uses()) {
  WorkList.push_back({VarUse.getUser()});
}
while (!WorkList.empty()) {
  auto &&WorkItem = WorkList.pop_back_val();
  auto *U = WorkItem.back();
  if (isa<llvm::ConstantExpr>(U)) {
    for (auto &&UU : U->uses()) {
      WorkItem.push_back(UU.getUser());
      WorkList.push_back(WorkItem);
      WorkItem.pop_back();
    }
    continue;
  }
  if (auto *I = dyn_cast<llvm::Instruction>(U)) {
    llvm::Value *OldV = Var;
    llvm::Instruction *NewV =
        new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false,
                           llvm::Align(Var->getAlignment()), I);
    WorkItem.pop_back();
    // Replace constant expressions directly or indirectly using the managed
    // variable with instructions.
    for (auto &&Op : WorkItem) {
      auto *CE = cast<llvm::ConstantExpr>(Op);
      auto *NewInst = llvm::createReplacementInstr(CE, I);
      NewInst->replaceUsesOfWith(OldV, NewV);
      OldV = CE;
      NewV = NewInst;
    }
    I->replaceUsesOfWith(OldV, NewV);
  } else {
    llvm_unreachable("Invalid use of managed variable")::llvm::llvm_unreachable_internal("Invalid use of managed variable"
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 475);
  }
}
478}

480/// Creates a function that sets up state on the host side for CUDA objects that
481/// have a presence on both the host and device sides. Specifically, registers
482/// the host side of kernel functions and device global variables with the CUDA
483/// runtime.
484/// \code
485/// void __cuda_register_globals(void** GpuBinaryHandle) {
486///    __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...);
487///    ...
488///    __cudaRegisterFunction(GpuBinaryHandle,KernelM,...);
489///    __cudaRegisterVar(GpuBinaryHandle, GlobalVar0, ...);
490///    ...
491///    __cudaRegisterVar(GpuBinaryHandle, GlobalVarN, ...);
492/// }
493/// \endcode
494llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
// No need to register anything
if (EmittedKernels.empty() && DeviceVars.empty())
  return nullptr;

llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
    getRegisterGlobalsFnTy(), llvm::GlobalValue::InternalLinkage,
    addUnderscoredPrefixToName("_register_globals"), &TheModule);
llvm::BasicBlock *EntryBB =
    llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc);
CGBuilderTy Builder(CGM, Context);
Builder.SetInsertPoint(EntryBB);

// void __cudaRegisterFunction(void **, const char *, char *, const char *,
//                             int, uint3*, uint3*, dim3*, dim3*, int*)
llvm::Type *RegisterFuncParams[] = {
    VoidPtrPtrTy, CharPtrTy, CharPtrTy, CharPtrTy, IntTy,
    VoidPtrTy,    VoidPtrTy, VoidPtrTy, VoidPtrTy, IntTy->getPointerTo()};
llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(IntTy, RegisterFuncParams, false),
    addUnderscoredPrefixToName("RegisterFunction"));

// Extract GpuBinaryHandle passed as the first argument passed to
// __cuda_register_globals() and generate __cudaRegisterFunction() call for
// each emitted kernel.
llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin();
for (auto &&I : EmittedKernels) {
  llvm::Constant *KernelName =
      makeConstantString(getDeviceSideName(cast<NamedDecl>(I.D)));
  llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(VoidPtrTy);
  llvm::Value *Args[] = {
      &GpuBinaryHandlePtr,
      Builder.CreateBitCast(KernelHandles[I.Kernel], VoidPtrTy),
      KernelName,
      KernelName,
      llvm::ConstantInt::get(IntTy, -1),
      NullPtr,
      NullPtr,
      NullPtr,
      NullPtr,
      llvm::ConstantPointerNull::get(IntTy->getPointerTo())};
  Builder.CreateCall(RegisterFunc, Args);
}

llvm::Type *VarSizeTy = IntTy;
// For HIP or CUDA 9.0+, device variable size is type of `size_t`.
if (CGM.getLangOpts().HIP ||
    ToCudaVersion(CGM.getTarget().getSDKVersion()) >= CudaVersion::CUDA_90)
  VarSizeTy = SizeTy;

// void __cudaRegisterVar(void **, char *, char *, const char *,
//                        int, int, int, int)
llvm::Type *RegisterVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
                                   CharPtrTy,    IntTy,     VarSizeTy,
                                   IntTy,        IntTy};
llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(VoidTy, RegisterVarParams, false),
    addUnderscoredPrefixToName("RegisterVar"));
// void __hipRegisterManagedVar(void **, char *, char *, const char *,
//                              size_t, unsigned)
llvm::Type *RegisterManagedVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
                                          CharPtrTy,    VarSizeTy, IntTy};
llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false),
    addUnderscoredPrefixToName("RegisterManagedVar"));
// void __cudaRegisterSurface(void **, const struct surfaceReference *,
//                            const void **, const char *, int, int);
llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(
        VoidTy, {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy},
        false),
    addUnderscoredPrefixToName("RegisterSurface"));
// void __cudaRegisterTexture(void **, const struct textureReference *,
//                            const void **, const char *, int, int, int)
llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(
        VoidTy,
        {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy, IntTy},
        false),
    addUnderscoredPrefixToName("RegisterTexture"));
for (auto &&Info : DeviceVars) {
  llvm::GlobalVariable *Var = Info.Var;
  assert((!Var->isDeclaration() || Info.Flags.isManaged()) &&(((!Var->isDeclaration() || Info.Flags.isManaged()) &&
 "External variables should not show up here, except HIP managed "
 "variables") ? static_cast<void> (0) : __assert_fail (
"(!Var->isDeclaration() || Info.Flags.isManaged()) && \"External variables should not show up here, except HIP managed \" \"variables\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 578, __PRETTY_FUNCTION__))
         "External variables should not show up here, except HIP managed "(((!Var->isDeclaration() || Info.Flags.isManaged()) &&
 "External variables should not show up here, except HIP managed "
 "variables") ? static_cast<void> (0) : __assert_fail (
"(!Var->isDeclaration() || Info.Flags.isManaged()) && \"External variables should not show up here, except HIP managed \" \"variables\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 578, __PRETTY_FUNCTION__))
         "variables")(((!Var->isDeclaration() || Info.Flags.isManaged()) &&
 "External variables should not show up here, except HIP managed "
 "variables") ? static_cast<void> (0) : __assert_fail (
"(!Var->isDeclaration() || Info.Flags.isManaged()) && \"External variables should not show up here, except HIP managed \" \"variables\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 578, __PRETTY_FUNCTION__));
  llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D));
  switch (Info.Flags.getKind()) {
  case DeviceVarFlags::Variable: {
    uint64_t VarSize =
        CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
    if (Info.Flags.isManaged()) {
      auto ManagedVar = new llvm::GlobalVariable(
          CGM.getModule(), Var->getType(),
          /*isConstant=*/false, Var->getLinkage(),
          /*Init=*/Var->isDeclaration()
              ? nullptr
              : llvm::ConstantPointerNull::get(Var->getType()),
          /*Name=*/"", /*InsertBefore=*/nullptr,
          llvm::GlobalVariable::NotThreadLocal);
      ManagedVar->setDSOLocal(Var->isDSOLocal());
      ManagedVar->setVisibility(Var->getVisibility());
      ManagedVar->setExternallyInitialized(true);
      ManagedVar->takeName(Var);
      Var->setName(Twine(ManagedVar->getName() + ".managed"));
      replaceManagedVar(Var, ManagedVar);
      llvm::Value *Args[] = {
          &GpuBinaryHandlePtr,
          Builder.CreateBitCast(ManagedVar, VoidPtrTy),
          Builder.CreateBitCast(Var, VoidPtrTy),
          VarName,
          llvm::ConstantInt::get(VarSizeTy, VarSize),
          llvm::ConstantInt::get(IntTy, Var->getAlignment())};
      if (!Var->isDeclaration())
        Builder.CreateCall(RegisterManagedVar, Args);
    } else {
      llvm::Value *Args[] = {
          &GpuBinaryHandlePtr,
          Builder.CreateBitCast(Var, VoidPtrTy),
          VarName,
          VarName,
          llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()),
          llvm::ConstantInt::get(VarSizeTy, VarSize),
          llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()),
          llvm::ConstantInt::get(IntTy, 0)};
      Builder.CreateCall(RegisterVar, Args);
    }
    break;
  }
  case DeviceVarFlags::Surface:
    Builder.CreateCall(
        RegisterSurf,
        {&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
         VarName, llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()),
         llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())});
    break;
  case DeviceVarFlags::Texture:
    Builder.CreateCall(
        RegisterTex,
        {&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
         VarName, llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()),
         llvm::ConstantInt::get(IntTy, Info.Flags.isNormalized()),
         llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())});
    break;
  }
}

Builder.CreateRetVoid();
return RegisterKernelsFunc;
642}

644/// Creates a global constructor function for the module:
645///
646/// For CUDA:
647/// \code
648/// void __cuda_module_ctor(void*) {
649///     Handle = __cudaRegisterFatBinary(GpuBinaryBlob);
650///     __cuda_register_globals(Handle);
651/// }
652/// \endcode
653///
654/// For HIP:
655/// \code
656/// void __hip_module_ctor(void*) {
657///     if (__hip_gpubin_handle == 0) {
658///         __hip_gpubin_handle  = __hipRegisterFatBinary(GpuBinaryBlob);
659///         __hip_register_globals(__hip_gpubin_handle);
660///     }
661/// }
662/// \endcode
663llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
bool IsHIP = CGM.getLangOpts().HIP;
bool IsCUDA = CGM.getLangOpts().CUDA;
// No need to generate ctors/dtors if there is no GPU binary.
StringRef CudaGpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName;
if (CudaGpuBinaryFileName.empty() && !IsHIP)
  return nullptr;
if ((IsHIP || (IsCUDA && !RelocatableDeviceCode)) && EmittedKernels.empty() &&
    DeviceVars.empty())
  return nullptr;

// void __{cuda|hip}_register_globals(void* handle);
llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
// We always need a function to pass in as callback. Create a dummy
// implementation if we don't need to register anything.
if (RelocatableDeviceCode && !RegisterGlobalsFunc)
  RegisterGlobalsFunc = makeDummyFunction(getRegisterGlobalsFnTy());

// void ** __{cuda|hip}RegisterFatBinary(void *);
llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(VoidPtrPtrTy, VoidPtrTy, false),
    addUnderscoredPrefixToName("RegisterFatBinary"));
// struct { int magic, int version, void * gpu_binary, void * dont_care };
llvm::StructType *FatbinWrapperTy =
    llvm::StructType::get(IntTy, IntTy, VoidPtrTy, VoidPtrTy);

// Register GPU binary with the CUDA runtime, store returned handle in a
// global variable and save a reference in GpuBinaryHandle to be cleaned up
// in destructor on exit. Then associate all known kernels with the GPU binary
// handle so CUDA runtime can figure out what to call on the GPU side.
std::unique_ptr<llvm::MemoryBuffer> CudaGpuBinary = nullptr;
if (!CudaGpuBinaryFileName.empty()) {
  llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> CudaGpuBinaryOrErr =
      llvm::MemoryBuffer::getFileOrSTDIN(CudaGpuBinaryFileName);
  if (std::error_code EC = CudaGpuBinaryOrErr.getError()) {
    CGM.getDiags().Report(diag::err_cannot_open_file)
        << CudaGpuBinaryFileName << EC.message();
    return nullptr;
  }
  CudaGpuBinary = std::move(CudaGpuBinaryOrErr.get());
}

llvm::Function *ModuleCtorFunc = llvm::Function::Create(
    llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
    llvm::GlobalValue::InternalLinkage,
    addUnderscoredPrefixToName("_module_ctor"), &TheModule);
llvm::BasicBlock *CtorEntryBB =
    llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc);
CGBuilderTy CtorBuilder(CGM, Context);

CtorBuilder.SetInsertPoint(CtorEntryBB);

const char *FatbinConstantName;
const char *FatbinSectionName;
const char *ModuleIDSectionName;
StringRef ModuleIDPrefix;
llvm::Constant *FatBinStr;
unsigned FatMagic;
if (IsHIP) {
  FatbinConstantName = ".hip_fatbin";
  FatbinSectionName = ".hipFatBinSegment";

  ModuleIDSectionName = "__hip_module_id";
  ModuleIDPrefix = "__hip_";

  if (CudaGpuBinary) {
    // If fatbin is available from early finalization, create a string
    // literal containing the fat binary loaded from the given file.
    const unsigned HIPCodeObjectAlign = 4096;
    FatBinStr =
        makeConstantString(std::string(CudaGpuBinary->getBuffer()), "",
                           FatbinConstantName, HIPCodeObjectAlign);
  } else {
    // If fatbin is not available, create an external symbol
    // __hip_fatbin in section .hip_fatbin. The external symbol is supposed
    // to contain the fat binary but will be populated somewhere else,
    // e.g. by lld through link script.
    FatBinStr = new llvm::GlobalVariable(
      CGM.getModule(), CGM.Int8Ty,
      /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr,
      "__hip_fatbin", nullptr,
      llvm::GlobalVariable::NotThreadLocal);
    cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName);
  }

  FatMagic = HIPFatMagic;
} else {
  if (RelocatableDeviceCode)
    FatbinConstantName = CGM.getTriple().isMacOSX()
                             ? "__NV_CUDA,__nv_relfatbin"
                             : "__nv_relfatbin";
  else
    FatbinConstantName =
        CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
  // NVIDIA's cuobjdump looks for fatbins in this section.
  FatbinSectionName =
      CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";

  ModuleIDSectionName = CGM.getTriple().isMacOSX()
                            ? "__NV_CUDA,__nv_module_id"
                            : "__nv_module_id";
  ModuleIDPrefix = "__nv_";

  // For CUDA, create a string literal containing the fat binary loaded from
  // the given file.
  FatBinStr = makeConstantString(std::string(CudaGpuBinary->getBuffer()), "",
                                 FatbinConstantName, 8);
  FatMagic = CudaFatMagic;
}

// Create initialized wrapper structure that points to the loaded GPU binary
ConstantInitBuilder Builder(CGM);
auto Values = Builder.beginStruct(FatbinWrapperTy);
// Fatbin wrapper magic.
Values.addInt(IntTy, FatMagic);
// Fatbin version.
Values.addInt(IntTy, 1);
// Data.
Values.add(FatBinStr);
// Unused in fatbin v1.
Values.add(llvm::ConstantPointerNull::get(VoidPtrTy));
llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
    addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
    /*constant*/ true);
FatbinWrapper->setSection(FatbinSectionName);

// There is only one HIP fat binary per linked module, however there are
// multiple constructor functions. Make sure the fat binary is registered
// only once. The constructor functions are executed by the dynamic loader
// before the program gains control. The dynamic loader cannot execute the
// constructor functions concurrently since doing that would not guarantee
// thread safety of the loaded program. Therefore we can assume sequential
// execution of constructor functions here.
if (IsHIP) {
  auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage :
      llvm::GlobalValue::LinkOnceAnyLinkage;
  llvm::BasicBlock *IfBlock =
      llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc);
  llvm::BasicBlock *ExitBlock =
      llvm::BasicBlock::Create(Context, "exit", ModuleCtorFunc);
  // The name, size, and initialization pattern of this variable is part
  // of HIP ABI.
  GpuBinaryHandle = new llvm::GlobalVariable(
      TheModule, VoidPtrPtrTy, /*isConstant=*/false,
      Linkage,
      /*Initializer=*/llvm::ConstantPointerNull::get(VoidPtrPtrTy),
      "__hip_gpubin_handle");
  GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
  // Prevent the weak symbol in different shared libraries being merged.
  if (Linkage != llvm::GlobalValue::InternalLinkage)
    GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
  Address GpuBinaryAddr(
      GpuBinaryHandle,
      CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
  {
    auto HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
    llvm::Constant *Zero =
        llvm::Constant::getNullValue(HandleValue->getType());
    llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(HandleValue, Zero);
    CtorBuilder.CreateCondBr(EQZero, IfBlock, ExitBlock);
  }
  {
    CtorBuilder.SetInsertPoint(IfBlock);
    // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper);
    llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
        RegisterFatbinFunc,
        CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
    CtorBuilder.CreateStore(RegisterFatbinCall, GpuBinaryAddr);
    CtorBuilder.CreateBr(ExitBlock);
  }
  {
    CtorBuilder.SetInsertPoint(ExitBlock);
    // Call __hip_register_globals(GpuBinaryHandle);
    if (RegisterGlobalsFunc) {
      auto HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
      CtorBuilder.CreateCall(RegisterGlobalsFunc, HandleValue);
    }
  }
} else if (!RelocatableDeviceCode) {
  // Register binary with CUDA runtime. This is substantially different in
  // default mode vs. separate compilation!
  // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
  llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
      RegisterFatbinFunc,
      CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
  GpuBinaryHandle = new llvm::GlobalVariable(
      TheModule, VoidPtrPtrTy, false, llvm::GlobalValue::InternalLinkage,
      llvm::ConstantPointerNull::get(VoidPtrPtrTy), "__cuda_gpubin_handle");
  GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
  CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
                                 CGM.getPointerAlign());

  // Call __cuda_register_globals(GpuBinaryHandle);
  if (RegisterGlobalsFunc)
    CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);

  // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it.
  if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
                         CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
    // void __cudaRegisterFatBinaryEnd(void **);
    llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
        llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
        "__cudaRegisterFatBinaryEnd");
    CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall);
  }
} else {
  // Generate a unique module ID.
  SmallString<64> ModuleID;
  llvm::raw_svector_ostream OS(ModuleID);
  OS << ModuleIDPrefix << llvm::format("%" PRIx64"l" "x", FatbinWrapper->getGUID());
  llvm::Constant *ModuleIDConstant = makeConstantString(
      std::string(ModuleID.str()), "", ModuleIDSectionName, 32);

  // Create an alias for the FatbinWrapper that nvcc will look for.
  llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage,
                            Twine("__fatbinwrap") + ModuleID, FatbinWrapper);

  // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *,
  // void *, void (*)(void **))
  SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
  RegisterLinkedBinaryName += ModuleID;
  llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
      getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName);

  assert(RegisterGlobalsFunc && "Expecting at least dummy function!")((RegisterGlobalsFunc && "Expecting at least dummy function!"
) ? static_cast<void> (0) : __assert_fail ("RegisterGlobalsFunc && \"Expecting at least dummy function!\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 887, __PRETTY_FUNCTION__));
  llvm::Value *Args[] = {RegisterGlobalsFunc,
                         CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy),
                         ModuleIDConstant,
                         makeDummyFunction(getCallbackFnTy())};
  CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args);
}

// Create destructor and register it with atexit() the way NVCC does it. Doing
// it during regular destructor phase worked in CUDA before 9.2 but results in
// double-free in 9.2.
if (llvm::Function *CleanupFn = makeModuleDtorFunction()) {
  // extern "C" int atexit(void (*f)(void));
  llvm::FunctionType *AtExitTy =
      llvm::FunctionType::get(IntTy, CleanupFn->getType(), false);
  llvm::FunctionCallee AtExitFunc =
      CGM.CreateRuntimeFunction(AtExitTy, "atexit", llvm::AttributeList(),
                                /*Local=*/true);
  CtorBuilder.CreateCall(AtExitFunc, CleanupFn);
}

CtorBuilder.CreateRetVoid();
return ModuleCtorFunc;
910}

912/// Creates a global destructor function that unregisters the GPU code blob
913/// registered by constructor.
914///
915/// For CUDA:
916/// \code
917/// void __cuda_module_dtor(void*) {
918///     __cudaUnregisterFatBinary(Handle);
919/// }
920/// \endcode
921///
922/// For HIP:
923/// \code
924/// void __hip_module_dtor(void*) {
925///     if (__hip_gpubin_handle) {
926///         __hipUnregisterFatBinary(__hip_gpubin_handle);
927///         __hip_gpubin_handle = 0;
928///     }
929/// }
930/// \endcode
931llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
// No need for destructor if we don't have a handle to unregister.
if (!GpuBinaryHandle)
  return nullptr;

// void __cudaUnregisterFatBinary(void ** handle);
llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
    llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
    addUnderscoredPrefixToName("UnregisterFatBinary"));

llvm::Function *ModuleDtorFunc = llvm::Function::Create(
    llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
    llvm::GlobalValue::InternalLinkage,
    addUnderscoredPrefixToName("_module_dtor"), &TheModule);

llvm::BasicBlock *DtorEntryBB =
    llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc);
CGBuilderTy DtorBuilder(CGM, Context);
DtorBuilder.SetInsertPoint(DtorEntryBB);

Address GpuBinaryAddr(GpuBinaryHandle, CharUnits::fromQuantity(
                                           GpuBinaryHandle->getAlignment()));
auto HandleValue = DtorBuilder.CreateLoad(GpuBinaryAddr);
// There is only one HIP fat binary per linked module, however there are
// multiple destructor functions. Make sure the fat binary is unregistered
// only once.
if (CGM.getLangOpts().HIP) {
  llvm::BasicBlock *IfBlock =
      llvm::BasicBlock::Create(Context, "if", ModuleDtorFunc);
  llvm::BasicBlock *ExitBlock =
      llvm::BasicBlock::Create(Context, "exit", ModuleDtorFunc);
  llvm::Constant *Zero = llvm::Constant::getNullValue(HandleValue->getType());
  llvm::Value *NEZero = DtorBuilder.CreateICmpNE(HandleValue, Zero);
  DtorBuilder.CreateCondBr(NEZero, IfBlock, ExitBlock);

  DtorBuilder.SetInsertPoint(IfBlock);
  DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
  DtorBuilder.CreateStore(Zero, GpuBinaryAddr);
  DtorBuilder.CreateBr(ExitBlock);

  DtorBuilder.SetInsertPoint(ExitBlock);
} else {
  DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
}
DtorBuilder.CreateRetVoid();
return ModuleDtorFunc;
977}

979CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
return new CGNVCUDARuntime(CGM);
981}

983void CGNVCUDARuntime::internalizeDeviceSideVar(
  const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) {
// For -fno-gpu-rdc, host-side shadows of external declarations of device-side
// global variables become internal definitions. These have to be internal in
// order to prevent name conflicts with global host variables with the same
// name in a different TUs.
//
// For -fgpu-rdc, the shadow variables should not be internalized because
// they may be accessed by different TU.
if (CGM.getLangOpts().GPURelocatableDeviceCode)
  return;

// __shared__ variables are odd. Shadows do get created, but
// they are not registered with the CUDA runtime, so they
// can't really be used to access their device-side
// counterparts. It's not clear yet whether it's nvcc's bug or
// a feature, but we've got to do the same for compatibility.
if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
    D->hasAttr<CUDASharedAttr>() ||
    D->getType()->isCUDADeviceBuiltinSurfaceType() ||
    D->getType()->isCUDADeviceBuiltinTextureType()) {
  Linkage = llvm::GlobalValue::InternalLinkage;
}
1006}

1008void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D,
                                          llvm::GlobalVariable &GV) {
if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
1
Calling 'Decl::hasAttr'→
4
←
Returning from 'Decl::hasAttr'→
5
←
Calling 'Decl::hasAttr'→
7
←
Returning from 'Decl::hasAttr'→
8
←
Taking false branch→
  // Shadow variables and their properties must be registered with CUDA
  // runtime. Skip Extern global variables, which will be registered in
  // the TU where they are defined.
  //
  // Don't register a C++17 inline variable. The local symbol can be
  // discarded and referencing a discarded local symbol from outside the
  // comdat (__cuda_register_globals) is disallowed by the ELF spec.
  // TODO: Reject __device__ constexpr and __device__ inline in Sema.
  // HIP managed variables need to be always recorded in device and host
  // compilations for transformation.
  if ((!D->hasExternalStorage() && !D->isInline()) ||
      D->hasAttr<HIPManagedAttr>()) {
    registerDeviceVar(D, GV, !D->hasDefinition(),
                      D->hasAttr<CUDAConstantAttr>());
  }
} else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
9
←
Assuming the condition is true→
10
←
Taking true branch→
           D->getType()->isCUDADeviceBuiltinTextureType()) {
  // Builtin surfaces and textures and their template arguments are
  // also registered with CUDA runtime.
  const ClassTemplateSpecializationDecl *TD =
      cast<ClassTemplateSpecializationDecl>(
          D->getType()->getAs<RecordType>()->getDecl());
11
←
Assuming the object is not a 'RecordType'→
12
←
Called C++ object pointer is null
  const TemplateArgumentList &Args = TD->getTemplateArgs();
  if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
    assert(Args.size() == 2 &&((Args.size() == 2 && "Unexpected number of template arguments of CUDA device "
 "builtin surface type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 2 && \"Unexpected number of template arguments of CUDA device \" \"builtin surface type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1037, __PRETTY_FUNCTION__))
           "Unexpected number of template arguments of CUDA device "((Args.size() == 2 && "Unexpected number of template arguments of CUDA device "
 "builtin surface type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 2 && \"Unexpected number of template arguments of CUDA device \" \"builtin surface type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1037, __PRETTY_FUNCTION__))
           "builtin surface type.")((Args.size() == 2 && "Unexpected number of template arguments of CUDA device "
 "builtin surface type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 2 && \"Unexpected number of template arguments of CUDA device \" \"builtin surface type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1037, __PRETTY_FUNCTION__));
    auto SurfType = Args[1].getAsIntegral();
    if (!D->hasExternalStorage())
      registerDeviceSurf(D, GV, !D->hasDefinition(), SurfType.getSExtValue());
  } else {
    assert(Args.size() == 3 &&((Args.size() == 3 && "Unexpected number of template arguments of CUDA device "
 "builtin texture type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 3 && \"Unexpected number of template arguments of CUDA device \" \"builtin texture type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1044, __PRETTY_FUNCTION__))
           "Unexpected number of template arguments of CUDA device "((Args.size() == 3 && "Unexpected number of template arguments of CUDA device "
 "builtin texture type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 3 && \"Unexpected number of template arguments of CUDA device \" \"builtin texture type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1044, __PRETTY_FUNCTION__))
           "builtin texture type.")((Args.size() == 3 && "Unexpected number of template arguments of CUDA device "
 "builtin texture type.") ? static_cast<void> (0) : __assert_fail
 ("Args.size() == 3 && \"Unexpected number of template arguments of CUDA device \" \"builtin texture type.\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1044, __PRETTY_FUNCTION__));
    auto TexType = Args[1].getAsIntegral();
    auto Normalized = Args[2].getAsIntegral();
    if (!D->hasExternalStorage())
      registerDeviceTex(D, GV, !D->hasDefinition(), TexType.getSExtValue(),
                        Normalized.getZExtValue());
  }
}
1052}

1054// Transform managed variables to pointers to managed variables in device code.
1055// Each use of the original managed variable is replaced by a load from the
1056// transformed managed variable. The transformed managed variable contains
1057// the address of managed memory which will be allocated by the runtime.
1058void CGNVCUDARuntime::transformManagedVars() {
for (auto &&Info : DeviceVars) {
  llvm::GlobalVariable *Var = Info.Var;
  if (Info.Flags.getKind() == DeviceVarFlags::Variable &&
      Info.Flags.isManaged()) {
    auto ManagedVar = new llvm::GlobalVariable(
        CGM.getModule(), Var->getType(),
        /*isConstant=*/false, Var->getLinkage(),
        /*Init=*/Var->isDeclaration()
            ? nullptr
            : llvm::ConstantPointerNull::get(Var->getType()),
        /*Name=*/"", /*InsertBefore=*/nullptr,
        llvm::GlobalVariable::NotThreadLocal,
        CGM.getContext().getTargetAddressSpace(LangAS::cuda_device));
    ManagedVar->setDSOLocal(Var->isDSOLocal());
    ManagedVar->setVisibility(Var->getVisibility());
    ManagedVar->setExternallyInitialized(true);
    replaceManagedVar(Var, ManagedVar);
    ManagedVar->takeName(Var);
    Var->setName(Twine(ManagedVar->getName()) + ".managed");
    // Keep managed variables even if they are not used in device code since
    // they need to be allocated by the runtime.
    if (!Var->isDeclaration()) {
      assert(!ManagedVar->isDeclaration())((!ManagedVar->isDeclaration()) ? static_cast<void> (
0) : __assert_fail ("!ManagedVar->isDeclaration()", "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/lib/CodeGen/CGCUDANV.cpp"
, 1081, __PRETTY_FUNCTION__));
      CGM.addCompilerUsedGlobal(Var);
      CGM.addCompilerUsedGlobal(ManagedVar);
    }
  }
}
1087}

1089// Returns module constructor to be added.
1090llvm::Function *CGNVCUDARuntime::finalizeModule() {
if (CGM.getLangOpts().CUDAIsDevice) {
  transformManagedVars();
  return nullptr;
}
return makeModuleCtorFunction();
1096}

1098llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
                                                  GlobalDecl GD) {
auto Loc = KernelHandles.find(F);
if (Loc != KernelHandles.end())
  return Loc->second;

if (!CGM.getLangOpts().HIP) {
  KernelHandles[F] = F;
  KernelStubs[F] = F;
  return F;
}

auto *Var = new llvm::GlobalVariable(
    TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(),
    /*Initializer=*/nullptr,
    CGM.getMangledName(
        GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel)));
Var->setAlignment(CGM.getPointerAlign().getAsAlign());
Var->setDSOLocal(F->isDSOLocal());
Var->setVisibility(F->getVisibility());
KernelHandles[F] = Var;
KernelStubs[Var] = F;
return Var;
1121}

←

/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h

1//===- DeclBase.h - Base Classes for representing declarations --*- C++ -*-===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9//  This file defines the Decl and DeclContext interfaces.
10//
11//===----------------------------------------------------------------------===//

13#ifndef LLVM_CLANG_AST_DECLBASE_H
14#define LLVM_CLANG_AST_DECLBASE_H

16#include "clang/AST/ASTDumperUtils.h"
17#include "clang/AST/AttrIterator.h"
18#include "clang/AST/DeclarationName.h"
19#include "clang/Basic/IdentifierTable.h"
20#include "clang/Basic/LLVM.h"
21#include "clang/Basic/SourceLocation.h"
22#include "clang/Basic/Specifiers.h"
23#include "llvm/ADT/ArrayRef.h"
24#include "llvm/ADT/PointerIntPair.h"
25#include "llvm/ADT/PointerUnion.h"
26#include "llvm/ADT/iterator.h"
27#include "llvm/ADT/iterator_range.h"
28#include "llvm/Support/Casting.h"
29#include "llvm/Support/Compiler.h"
30#include "llvm/Support/PrettyStackTrace.h"
31#include "llvm/Support/VersionTuple.h"
32#include <algorithm>
33#include <cassert>
34#include <cstddef>
35#include <iterator>
36#include <string>
37#include <type_traits>
38#include <utility>

40namespace clang {

42class ASTContext;
43class ASTMutationListener;
44class Attr;
45class BlockDecl;
46class DeclContext;
47class ExternalSourceSymbolAttr;
48class FunctionDecl;
49class FunctionType;
50class IdentifierInfo;
51enum Linkage : unsigned char;
52class LinkageSpecDecl;
53class Module;
54class NamedDecl;
55class ObjCCategoryDecl;
56class ObjCCategoryImplDecl;
57class ObjCContainerDecl;
58class ObjCImplDecl;
59class ObjCImplementationDecl;
60class ObjCInterfaceDecl;
61class ObjCMethodDecl;
62class ObjCProtocolDecl;
63struct PrintingPolicy;
64class RecordDecl;
65class SourceManager;
66class Stmt;
67class StoredDeclsMap;
68class TemplateDecl;
69class TemplateParameterList;
70class TranslationUnitDecl;
71class UsingDirectiveDecl;

73/// Captures the result of checking the availability of a
74/// declaration.
75enum AvailabilityResult {
AR_Available = 0,
AR_NotYetIntroduced,
AR_Deprecated,
AR_Unavailable
80};

82/// Decl - This represents one declaration (or definition), e.g. a variable,
83/// typedef, function, struct, etc.
84///
85/// Note: There are objects tacked on before the *beginning* of Decl
86/// (and its subclasses) in its Decl::operator new(). Proper alignment
87/// of all subclasses (not requiring more than the alignment of Decl) is
88/// asserted in DeclBase.cpp.
89class alignas(8) Decl {
90public:
/// Lists the kind of concrete classes of Decl.
enum Kind {
93#define DECL(DERIVED, BASE) DERIVED,
94#define ABSTRACT_DECL(DECL)
95#define DECL_RANGE(BASE, START, END) \
      first##BASE = START, last##BASE = END,
97#define LAST_DECL_RANGE(BASE, START, END) \
      first##BASE = START, last##BASE = END
99#include "clang/AST/DeclNodes.inc"
};

/// A placeholder type used to construct an empty shell of a
/// decl-derived type that will be filled in later (e.g., by some
/// deserialization method).
struct EmptyShell {};

/// IdentifierNamespace - The different namespaces in which
/// declarations may appear.  According to C99 6.2.3, there are
/// four namespaces, labels, tags, members and ordinary
/// identifiers.  C++ describes lookup completely differently:
/// certain lookups merely "ignore" certain kinds of declarations,
/// usually based on whether the declaration is of a type, etc.
///
/// These are meant as bitmasks, so that searches in
/// C++ can look into the "tag" namespace during ordinary lookup.
///
/// Decl currently provides 15 bits of IDNS bits.
enum IdentifierNamespace {
  /// Labels, declared with 'x:' and referenced with 'goto x'.
  IDNS_Label               = 0x0001,

  /// Tags, declared with 'struct foo;' and referenced with
  /// 'struct foo'.  All tags are also types.  This is what
  /// elaborated-type-specifiers look for in C.
  /// This also contains names that conflict with tags in the
  /// same scope but that are otherwise ordinary names (non-type
  /// template parameters and indirect field declarations).
  IDNS_Tag                 = 0x0002,

  /// Types, declared with 'struct foo', typedefs, etc.
  /// This is what elaborated-type-specifiers look for in C++,
  /// but note that it's ill-formed to find a non-tag.
  IDNS_Type                = 0x0004,

  /// Members, declared with object declarations within tag
  /// definitions.  In C, these can only be found by "qualified"
  /// lookup in member expressions.  In C++, they're found by
  /// normal lookup.
  IDNS_Member              = 0x0008,

  /// Namespaces, declared with 'namespace foo {}'.
  /// Lookup for nested-name-specifiers find these.
  IDNS_Namespace           = 0x0010,

  /// Ordinary names.  In C, everything that's not a label, tag,
  /// member, or function-local extern ends up here.
  IDNS_Ordinary            = 0x0020,

  /// Objective C \@protocol.
  IDNS_ObjCProtocol        = 0x0040,

  /// This declaration is a friend function.  A friend function
  /// declaration is always in this namespace but may also be in
  /// IDNS_Ordinary if it was previously declared.
  IDNS_OrdinaryFriend      = 0x0080,

  /// This declaration is a friend class.  A friend class
  /// declaration is always in this namespace but may also be in
  /// IDNS_Tag|IDNS_Type if it was previously declared.
  IDNS_TagFriend           = 0x0100,

  /// This declaration is a using declaration.  A using declaration
  /// *introduces* a number of other declarations into the current
  /// scope, and those declarations use the IDNS of their targets,
  /// but the actual using declarations go in this namespace.
  IDNS_Using               = 0x0200,

  /// This declaration is a C++ operator declared in a non-class
  /// context.  All such operators are also in IDNS_Ordinary.
  /// C++ lexical operator lookup looks for these.
  IDNS_NonMemberOperator   = 0x0400,

  /// This declaration is a function-local extern declaration of a
  /// variable or function. This may also be IDNS_Ordinary if it
  /// has been declared outside any function. These act mostly like
  /// invisible friend declarations, but are also visible to unqualified
  /// lookup within the scope of the declaring function.
  IDNS_LocalExtern         = 0x0800,

  /// This declaration is an OpenMP user defined reduction construction.
  IDNS_OMPReduction        = 0x1000,

  /// This declaration is an OpenMP user defined mapper.
  IDNS_OMPMapper           = 0x2000,
};

/// ObjCDeclQualifier - 'Qualifiers' written next to the return and
/// parameter types in method declarations.  Other than remembering
/// them and mangling them into the method's signature string, these
/// are ignored by the compiler; they are consumed by certain
/// remote-messaging frameworks.
///
/// in, inout, and out are mutually exclusive and apply only to
/// method parameters.  bycopy and byref are mutually exclusive and
/// apply only to method parameters (?).  oneway applies only to
/// results.  All of these expect their corresponding parameter to
/// have a particular type.  None of this is currently enforced by
/// clang.
///
/// This should be kept in sync with ObjCDeclSpec::ObjCDeclQualifier.
enum ObjCDeclQualifier {
  OBJC_TQ_None = 0x0,
  OBJC_TQ_In = 0x1,
  OBJC_TQ_Inout = 0x2,
  OBJC_TQ_Out = 0x4,
  OBJC_TQ_Bycopy = 0x8,
  OBJC_TQ_Byref = 0x10,
  OBJC_TQ_Oneway = 0x20,

  /// The nullability qualifier is set when the nullability of the
  /// result or parameter was expressed via a context-sensitive
  /// keyword.
  OBJC_TQ_CSNullability = 0x40
};

/// The kind of ownership a declaration has, for visibility purposes.
/// This enumeration is designed such that higher values represent higher
/// levels of name hiding.
enum class ModuleOwnershipKind : unsigned {
  /// This declaration is not owned by a module.
  Unowned,

  /// This declaration has an owning module, but is globally visible
  /// (typically because its owning module is visible and we know that
  /// modules cannot later become hidden in this compilation).
  /// After serialization and deserialization, this will be converted
  /// to VisibleWhenImported.
  Visible,

  /// This declaration has an owning module, and is visible when that
  /// module is imported.
  VisibleWhenImported,

  /// This declaration has an owning module, but is only visible to
  /// lookups that occur within that module.
  ModulePrivate
};

239protected:
/// The next declaration within the same lexical
/// DeclContext. These pointers form the linked list that is
/// traversed via DeclContext's decls_begin()/decls_end().
///
/// The extra two bits are used for the ModuleOwnershipKind.
llvm::PointerIntPair<Decl *, 2, ModuleOwnershipKind> NextInContextAndBits;

247private:
friend class DeclContext;

struct MultipleDC {
  DeclContext *SemanticDC;
  DeclContext *LexicalDC;
};

/// DeclCtx - Holds either a DeclContext* or a MultipleDC*.
/// For declarations that don't contain C++ scope specifiers, it contains
/// the DeclContext where the Decl was declared.
/// For declarations with C++ scope specifiers, it contains a MultipleDC*
/// with the context where it semantically belongs (SemanticDC) and the
/// context where it was lexically declared (LexicalDC).
/// e.g.:
///
///   namespace A {
///      void f(); // SemanticDC == LexicalDC == 'namespace A'
///   }
///   void A::f(); // SemanticDC == namespace 'A'
///                // LexicalDC == global namespace
llvm::PointerUnion<DeclContext*, MultipleDC*> DeclCtx;

bool isInSemaDC() const { return DeclCtx.is<DeclContext*>(); }
bool isOutOfSemaDC() const { return DeclCtx.is<MultipleDC*>(); }

MultipleDC *getMultipleDC() const {
  return DeclCtx.get<MultipleDC*>();
}

DeclContext *getSemanticDC() const {
  return DeclCtx.get<DeclContext*>();
}

/// Loc - The location of this decl.
SourceLocation Loc;

/// DeclKind - This indicates which class this is.
unsigned DeclKind : 7;

/// InvalidDecl - This indicates a semantic error occurred.
unsigned InvalidDecl :  1;

/// HasAttrs - This indicates whether the decl has attributes or not.
unsigned HasAttrs : 1;

/// Implicit - Whether this declaration was implicitly generated by
/// the implementation rather than explicitly written by the user.
unsigned Implicit : 1;

/// Whether this declaration was "used", meaning that a definition is
/// required.
unsigned Used : 1;

/// Whether this declaration was "referenced".
/// The difference with 'Used' is whether the reference appears in a
/// evaluated context or not, e.g. functions used in uninstantiated templates
/// are regarded as "referenced" but not "used".
unsigned Referenced : 1;

/// Whether this declaration is a top-level declaration (function,
/// global variable, etc.) that is lexically inside an objc container
/// definition.
unsigned TopLevelDeclInObjCContainer : 1;

/// Whether statistic collection is enabled.
static bool StatisticsEnabled;

315protected:
friend class ASTDeclReader;
friend class ASTDeclWriter;
friend class ASTNodeImporter;
friend class ASTReader;
friend class CXXClassMemberWrapper;
friend class LinkageComputer;
template<typename decl_type> friend class Redeclarable;

/// Access - Used by C++ decls for the access specifier.
// NOTE: VC++ treats enums as signed, avoid using the AccessSpecifier enum
unsigned Access : 2;

/// Whether this declaration was loaded from an AST file.
unsigned FromASTFile : 1;

/// IdentifierNamespace - This specifies what IDNS_* namespace this lives in.
unsigned IdentifierNamespace : 14;

/// If 0, we have not computed the linkage of this declaration.
/// Otherwise, it is the linkage + 1.
mutable unsigned CacheValidAndLinkage : 3;

/// Allocate memory for a deserialized declaration.
///
/// This routine must be used to allocate memory for any declaration that is
/// deserialized from a module file.
///
/// \param Size The size of the allocated object.
/// \param Ctx The context in which we will allocate memory.
/// \param ID The global ID of the deserialized declaration.
/// \param Extra The amount of extra space to allocate after the object.
void *operator new(std::size_t Size, const ASTContext &Ctx, unsigned ID,
                   std::size_t Extra = 0);

/// Allocate memory for a non-deserialized declaration.
void *operator new(std::size_t Size, const ASTContext &Ctx,
                   DeclContext *Parent, std::size_t Extra = 0);

354private:
bool AccessDeclContextSanity() const;

/// Get the module ownership kind to use for a local lexical child of \p DC,
/// which may be either a local or (rarely) an imported declaration.
static ModuleOwnershipKind getModuleOwnershipKindForChildOf(DeclContext *DC) {
  if (DC) {
    auto *D = cast<Decl>(DC);
    auto MOK = D->getModuleOwnershipKind();
    if (MOK != ModuleOwnershipKind::Unowned &&
        (!D->isFromASTFile() || D->hasLocalOwningModuleStorage()))
      return MOK;
    // If D is not local and we have no local module storage, then we don't
    // need to track module ownership at all.
  }
  return ModuleOwnershipKind::Unowned;
}

372public:
Decl() = delete;
Decl(const Decl&) = delete;
Decl(Decl &&) = delete;
Decl &operator=(const Decl&) = delete;
Decl &operator=(Decl&&) = delete;

379protected:
Decl(Kind DK, DeclContext *DC, SourceLocation L)
    : NextInContextAndBits(nullptr, getModuleOwnershipKindForChildOf(DC)),
      DeclCtx(DC), Loc(L), DeclKind(DK), InvalidDecl(false), HasAttrs(false),
      Implicit(false), Used(false), Referenced(false),
      TopLevelDeclInObjCContainer(false), Access(AS_none), FromASTFile(0),
      IdentifierNamespace(getIdentifierNamespaceForKind(DK)),
      CacheValidAndLinkage(0) {
  if (StatisticsEnabled) add(DK);
}

Decl(Kind DK, EmptyShell Empty)
    : DeclKind(DK), InvalidDecl(false), HasAttrs(false), Implicit(false),
      Used(false), Referenced(false), TopLevelDeclInObjCContainer(false),
      Access(AS_none), FromASTFile(0),
      IdentifierNamespace(getIdentifierNamespaceForKind(DK)),
      CacheValidAndLinkage(0) {
  if (StatisticsEnabled) add(DK);
}

virtual ~Decl();

/// Update a potentially out-of-date declaration.
void updateOutOfDate(IdentifierInfo &II) const;

Linkage getCachedLinkage() const {
  return Linkage(CacheValidAndLinkage - 1);
}

void setCachedLinkage(Linkage L) const {
  CacheValidAndLinkage = L + 1;
}

bool hasCachedLinkage() const {
  return CacheValidAndLinkage;
}

416public:
/// Source range that this declaration covers.
virtual SourceRange getSourceRange() const LLVM_READONLY__attribute__((__pure__)) {
  return SourceRange(getLocation(), getLocation());
}

SourceLocation getBeginLoc() const LLVM_READONLY__attribute__((__pure__)) {
  return getSourceRange().getBegin();
}

SourceLocation getEndLoc() const LLVM_READONLY__attribute__((__pure__)) {
  return getSourceRange().getEnd();
}

SourceLocation getLocation() const { return Loc; }
void setLocation(SourceLocation L) { Loc = L; }

Kind getKind() const { return static_cast<Kind>(DeclKind); }
const char *getDeclKindName() const;

Decl *getNextDeclInContext() { return NextInContextAndBits.getPointer(); }
const Decl *getNextDeclInContext() const {return NextInContextAndBits.getPointer();}

DeclContext *getDeclContext() {
  if (isInSemaDC())
    return getSemanticDC();
  return getMultipleDC()->SemanticDC;
}
const DeclContext *getDeclContext() const {
  return const_cast<Decl*>(this)->getDeclContext();
}

/// Find the innermost non-closure ancestor of this declaration,
/// walking up through blocks, lambdas, etc.  If that ancestor is
/// not a code context (!isFunctionOrMethod()), returns null.
///
/// A declaration may be its own non-closure context.
Decl *getNonClosureContext();
const Decl *getNonClosureContext() const {
  return const_cast<Decl*>(this)->getNonClosureContext();
}

TranslationUnitDecl *getTranslationUnitDecl();
const TranslationUnitDecl *getTranslationUnitDecl() const {
  return const_cast<Decl*>(this)->getTranslationUnitDecl();
}

bool isInAnonymousNamespace() const;

bool isInStdNamespace() const;

ASTContext &getASTContext() const LLVM_READONLY__attribute__((__pure__));

/// Helper to get the language options from the ASTContext.
/// Defined out of line to avoid depending on ASTContext.h.
const LangOptions &getLangOpts() const LLVM_READONLY__attribute__((__pure__));

void setAccess(AccessSpecifier AS) {
  Access = AS;
  assert(AccessDeclContextSanity())((AccessDeclContextSanity()) ? static_cast<void> (0) : __assert_fail
 ("AccessDeclContextSanity()", "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 475, __PRETTY_FUNCTION__));
}

AccessSpecifier getAccess() const {
  assert(AccessDeclContextSanity())((AccessDeclContextSanity()) ? static_cast<void> (0) : __assert_fail
 ("AccessDeclContextSanity()", "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 479, __PRETTY_FUNCTION__));
  return AccessSpecifier(Access);
}

/// Retrieve the access specifier for this declaration, even though
/// it may not yet have been properly set.
AccessSpecifier getAccessUnsafe() const {
  return AccessSpecifier(Access);
}

bool hasAttrs() const { return HasAttrs; }

void setAttrs(const AttrVec& Attrs) {
  return setAttrsImpl(Attrs, getASTContext());
}

AttrVec &getAttrs() {
  return const_cast<AttrVec&>(const_cast<const Decl*>(this)->getAttrs());
}

const AttrVec &getAttrs() const;
void dropAttrs();
void addAttr(Attr *A);

using attr_iterator = AttrVec::const_iterator;
using attr_range = llvm::iterator_range<attr_iterator>;

attr_range attrs() const {
  return attr_range(attr_begin(), attr_end());
}

attr_iterator attr_begin() const {
  return hasAttrs() ? getAttrs().begin() : nullptr;
}
attr_iterator attr_end() const {
  return hasAttrs() ? getAttrs().end() : nullptr;
}

template <typename T>
void dropAttr() {
  if (!HasAttrs) return;

  AttrVec &Vec = getAttrs();
  llvm::erase_if(Vec, [](Attr *A) { return isa<T>(A); });

  if (Vec.empty())
    HasAttrs = false;
}

template <typename T>
llvm::iterator_range<specific_attr_iterator<T>> specific_attrs() const {
  return llvm::make_range(specific_attr_begin<T>(), specific_attr_end<T>());
}

template <typename T>
specific_attr_iterator<T> specific_attr_begin() const {
  return specific_attr_iterator<T>(attr_begin());
}

template <typename T>
specific_attr_iterator<T> specific_attr_end() const {
  return specific_attr_iterator<T>(attr_end());
}

template<typename T> T *getAttr() const {
  return hasAttrs() ? getSpecificAttr<T>(getAttrs()) : nullptr;
}

template<typename T> bool hasAttr() const {
  return hasAttrs() && hasSpecificAttr<T>(getAttrs());
2
←
Assuming the condition is false→
3
←
Returning zero, which participates in a condition later→
6
←
Returning zero, which participates in a condition later→
}

/// getMaxAlignment - return the maximum alignment specified by attributes
/// on this decl, 0 if there are none.
unsigned getMaxAlignment() const;

/// setInvalidDecl - Indicates the Decl had a semantic error. This
/// allows for graceful error recovery.
void setInvalidDecl(bool Invalid = true);
bool isInvalidDecl() const { return (bool) InvalidDecl; }

/// isImplicit - Indicates whether the declaration was implicitly
/// generated by the implementation. If false, this declaration
/// was written explicitly in the source code.
bool isImplicit() const { return Implicit; }
void setImplicit(bool I = true) { Implicit = I; }

/// Whether *any* (re-)declaration of the entity was used, meaning that
/// a definition is required.
///
/// \param CheckUsedAttr When true, also consider the "used" attribute
/// (in addition to the "used" bit set by \c setUsed()) when determining
/// whether the function is used.
bool isUsed(bool CheckUsedAttr = true) const;

/// Set whether the declaration is used, in the sense of odr-use.
///
/// This should only be used immediately after creating a declaration.
/// It intentionally doesn't notify any listeners.
void setIsUsed() { getCanonicalDecl()->Used = true; }

/// Mark the declaration used, in the sense of odr-use.
///
/// This notifies any mutation listeners in addition to setting a bit
/// indicating the declaration is used.
void markUsed(ASTContext &C);

/// Whether any declaration of this entity was referenced.
bool isReferenced() const;

/// Whether this declaration was referenced. This should not be relied
/// upon for anything other than debugging.
bool isThisDeclarationReferenced() const { return Referenced; }

void setReferenced(bool R = true) { Referenced = R; }

/// Whether this declaration is a top-level declaration (function,
/// global variable, etc.) that is lexically inside an objc container
/// definition.
bool isTopLevelDeclInObjCContainer() const {
  return TopLevelDeclInObjCContainer;
}

void setTopLevelDeclInObjCContainer(bool V = true) {
  TopLevelDeclInObjCContainer = V;
}

/// Looks on this and related declarations for an applicable
/// external source symbol attribute.
ExternalSourceSymbolAttr *getExternalSourceSymbolAttr() const;

/// Whether this declaration was marked as being private to the
/// module in which it was defined.
bool isModulePrivate() const {
  return getModuleOwnershipKind() == ModuleOwnershipKind::ModulePrivate;
}

/// Return true if this declaration has an attribute which acts as
/// definition of the entity, such as 'alias' or 'ifunc'.
bool hasDefiningAttr() const;

/// Return this declaration's defining attribute if it has one.
const Attr *getDefiningAttr() const;

623protected:
/// Specify that this declaration was marked as being private
/// to the module in which it was defined.
void setModulePrivate() {
  // The module-private specifier has no effect on unowned declarations.
  // FIXME: We should track this in some way for source fidelity.
  if (getModuleOwnershipKind() == ModuleOwnershipKind::Unowned)
    return;
  setModuleOwnershipKind(ModuleOwnershipKind::ModulePrivate);
}

634public:
/// Set the FromASTFile flag. This indicates that this declaration
/// was deserialized and not parsed from source code and enables
/// features such as module ownership information.
void setFromASTFile() {
  FromASTFile = true;
}

/// Set the owning module ID.  This may only be called for
/// deserialized Decls.
void setOwningModuleID(unsigned ID) {
  assert(isFromASTFile() && "Only works on a deserialized declaration")((isFromASTFile() && "Only works on a deserialized declaration"
) ? static_cast<void> (0) : __assert_fail ("isFromASTFile() && \"Only works on a deserialized declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 645, __PRETTY_FUNCTION__));
  *((unsigned*)this - 2) = ID;
}

649public:
/// Determine the availability of the given declaration.
///
/// This routine will determine the most restrictive availability of
/// the given declaration (e.g., preferring 'unavailable' to
/// 'deprecated').
///
/// \param Message If non-NULL and the result is not \c
/// AR_Available, will be set to a (possibly empty) message
/// describing why the declaration has not been introduced, is
/// deprecated, or is unavailable.
///
/// \param EnclosingVersion The version to compare with. If empty, assume the
/// deployment target version.
///
/// \param RealizedPlatform If non-NULL and the availability result is found
/// in an available attribute it will set to the platform which is written in
/// the available attribute.
AvailabilityResult
getAvailability(std::string *Message = nullptr,
                VersionTuple EnclosingVersion = VersionTuple(),
                StringRef *RealizedPlatform = nullptr) const;

/// Retrieve the version of the target platform in which this
/// declaration was introduced.
///
/// \returns An empty version tuple if this declaration has no 'introduced'
/// availability attributes, or the version tuple that's specified in the
/// attribute otherwise.
VersionTuple getVersionIntroduced() const;

/// Determine whether this declaration is marked 'deprecated'.
///
/// \param Message If non-NULL and the declaration is deprecated,
/// this will be set to the message describing why the declaration
/// was deprecated (which may be empty).
bool isDeprecated(std::string *Message = nullptr) const {
  return getAvailability(Message) == AR_Deprecated;
}

/// Determine whether this declaration is marked 'unavailable'.
///
/// \param Message If non-NULL and the declaration is unavailable,
/// this will be set to the message describing why the declaration
/// was made unavailable (which may be empty).
bool isUnavailable(std::string *Message = nullptr) const {
  return getAvailability(Message) == AR_Unavailable;
}

/// Determine whether this is a weak-imported symbol.
///
/// Weak-imported symbols are typically marked with the
/// 'weak_import' attribute, but may also be marked with an
/// 'availability' attribute where we're targing a platform prior to
/// the introduction of this feature.
bool isWeakImported() const;

/// Determines whether this symbol can be weak-imported,
/// e.g., whether it would be well-formed to add the weak_import
/// attribute.
///
/// \param IsDefinition Set to \c true to indicate that this
/// declaration cannot be weak-imported because it has a definition.
bool canBeWeakImported(bool &IsDefinition) const;

/// Determine whether this declaration came from an AST file (such as
/// a precompiled header or module) rather than having been parsed.
bool isFromASTFile() const { return FromASTFile; }

/// Retrieve the global declaration ID associated with this
/// declaration, which specifies where this Decl was loaded from.
unsigned getGlobalID() const {
  if (isFromASTFile())
    return *((const unsigned*)this - 1);
  return 0;
}

/// Retrieve the global ID of the module that owns this particular
/// declaration.
unsigned getOwningModuleID() const {
  if (isFromASTFile())
    return *((const unsigned*)this - 2);
  return 0;
}

734private:
Module *getOwningModuleSlow() const;

737protected:
bool hasLocalOwningModuleStorage() const;

740public:
/// Get the imported owning module, if this decl is from an imported
/// (non-local) module.
Module *getImportedOwningModule() const {
  if (!isFromASTFile() || !hasOwningModule())
    return nullptr;

  return getOwningModuleSlow();
}

/// Get the local owning module, if known. Returns nullptr if owner is
/// not yet known or declaration is not from a module.
Module *getLocalOwningModule() const {
  if (isFromASTFile() || !hasOwningModule())
    return nullptr;

  assert(hasLocalOwningModuleStorage() &&((hasLocalOwningModuleStorage() && "owned local decl but no local module storage"
) ? static_cast<void> (0) : __assert_fail ("hasLocalOwningModuleStorage() && \"owned local decl but no local module storage\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 757, __PRETTY_FUNCTION__))
         "owned local decl but no local module storage")((hasLocalOwningModuleStorage() && "owned local decl but no local module storage"
) ? static_cast<void> (0) : __assert_fail ("hasLocalOwningModuleStorage() && \"owned local decl but no local module storage\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 757, __PRETTY_FUNCTION__));
  return reinterpret_cast<Module *const *>(this)[-1];
}
void setLocalOwningModule(Module *M) {
  assert(!isFromASTFile() && hasOwningModule() &&((!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage
() && "should not have a cached owning module") ? static_cast
<void> (0) : __assert_fail ("!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage() && \"should not have a cached owning module\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 763, __PRETTY_FUNCTION__))
         hasLocalOwningModuleStorage() &&((!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage
() && "should not have a cached owning module") ? static_cast
<void> (0) : __assert_fail ("!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage() && \"should not have a cached owning module\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 763, __PRETTY_FUNCTION__))
         "should not have a cached owning module")((!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage
() && "should not have a cached owning module") ? static_cast
<void> (0) : __assert_fail ("!isFromASTFile() && hasOwningModule() && hasLocalOwningModuleStorage() && \"should not have a cached owning module\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 763, __PRETTY_FUNCTION__));
  reinterpret_cast<Module **>(this)[-1] = M;
}

/// Is this declaration owned by some module?
bool hasOwningModule() const {
  return getModuleOwnershipKind() != ModuleOwnershipKind::Unowned;
}

/// Get the module that owns this declaration (for visibility purposes).
Module *getOwningModule() const {
  return isFromASTFile() ? getImportedOwningModule() : getLocalOwningModule();
}

/// Get the module that owns this declaration for linkage purposes.
/// There only ever is such a module under the C++ Modules TS.
///
/// \param IgnoreLinkage Ignore the linkage of the entity; assume that
/// all declarations in a global module fragment are unowned.
Module *getOwningModuleForLinkage(bool IgnoreLinkage = false) const;

/// Determine whether this declaration is definitely visible to name lookup,
/// independent of whether the owning module is visible.
/// Note: The declaration may be visible even if this returns \c false if the
/// owning module is visible within the query context. This is a low-level
/// helper function; most code should be calling Sema::isVisible() instead.
bool isUnconditionallyVisible() const {
  return (int)getModuleOwnershipKind() <= (int)ModuleOwnershipKind::Visible;
}

/// Set that this declaration is globally visible, even if it came from a
/// module that is not visible.
void setVisibleDespiteOwningModule() {
  if (!isUnconditionallyVisible())
    setModuleOwnershipKind(ModuleOwnershipKind::Visible);
}

/// Get the kind of module ownership for this declaration.
ModuleOwnershipKind getModuleOwnershipKind() const {
  return NextInContextAndBits.getInt();
}

/// Set whether this declaration is hidden from name lookup.
void setModuleOwnershipKind(ModuleOwnershipKind MOK) {
  assert(!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned &&((!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned &&
 MOK != ModuleOwnershipKind::Unowned && !isFromASTFile
() && !hasLocalOwningModuleStorage()) && "no storage available for owning module for this declaration"
) ? static_cast<void> (0) : __assert_fail ("!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned && MOK != ModuleOwnershipKind::Unowned && !isFromASTFile() && !hasLocalOwningModuleStorage()) && \"no storage available for owning module for this declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 810, __PRETTY_FUNCTION__))
           MOK != ModuleOwnershipKind::Unowned && !isFromASTFile() &&((!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned &&
 MOK != ModuleOwnershipKind::Unowned && !isFromASTFile
() && !hasLocalOwningModuleStorage()) && "no storage available for owning module for this declaration"
) ? static_cast<void> (0) : __assert_fail ("!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned && MOK != ModuleOwnershipKind::Unowned && !isFromASTFile() && !hasLocalOwningModuleStorage()) && \"no storage available for owning module for this declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 810, __PRETTY_FUNCTION__))
           !hasLocalOwningModuleStorage()) &&((!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned &&
 MOK != ModuleOwnershipKind::Unowned && !isFromASTFile
() && !hasLocalOwningModuleStorage()) && "no storage available for owning module for this declaration"
) ? static_cast<void> (0) : __assert_fail ("!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned && MOK != ModuleOwnershipKind::Unowned && !isFromASTFile() && !hasLocalOwningModuleStorage()) && \"no storage available for owning module for this declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 810, __PRETTY_FUNCTION__))
         "no storage available for owning module for this declaration")((!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned &&
 MOK != ModuleOwnershipKind::Unowned && !isFromASTFile
() && !hasLocalOwningModuleStorage()) && "no storage available for owning module for this declaration"
) ? static_cast<void> (0) : __assert_fail ("!(getModuleOwnershipKind() == ModuleOwnershipKind::Unowned && MOK != ModuleOwnershipKind::Unowned && !isFromASTFile() && !hasLocalOwningModuleStorage()) && \"no storage available for owning module for this declaration\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 810, __PRETTY_FUNCTION__));
  NextInContextAndBits.setInt(MOK);
}

unsigned getIdentifierNamespace() const {
  return IdentifierNamespace;
}

bool isInIdentifierNamespace(unsigned NS) const {
  return getIdentifierNamespace() & NS;
}

static unsigned getIdentifierNamespaceForKind(Kind DK);

bool hasTagIdentifierNamespace() const {
  return isTagIdentifierNamespace(getIdentifierNamespace());
}

static bool isTagIdentifierNamespace(unsigned NS) {
  // TagDecls have Tag and Type set and may also have TagFriend.
  return (NS & ~IDNS_TagFriend) == (IDNS_Tag | IDNS_Type);
}

/// getLexicalDeclContext - The declaration context where this Decl was
/// lexically declared (LexicalDC). May be different from
/// getDeclContext() (SemanticDC).
/// e.g.:
///
///   namespace A {
///      void f(); // SemanticDC == LexicalDC == 'namespace A'
///   }
///   void A::f(); // SemanticDC == namespace 'A'
///                // LexicalDC == global namespace
DeclContext *getLexicalDeclContext() {
  if (isInSemaDC())
    return getSemanticDC();
  return getMultipleDC()->LexicalDC;
}
const DeclContext *getLexicalDeclContext() const {
  return const_cast<Decl*>(this)->getLexicalDeclContext();
}

/// Determine whether this declaration is declared out of line (outside its
/// semantic context).
virtual bool isOutOfLine() const;

/// setDeclContext - Set both the semantic and lexical DeclContext
/// to DC.
void setDeclContext(DeclContext *DC);

void setLexicalDeclContext(DeclContext *DC);

/// Determine whether this declaration is a templated entity (whether it is
// within the scope of a template parameter).
bool isTemplated() const;

/// Determine the number of levels of template parameter surrounding this
/// declaration.
unsigned getTemplateDepth() const;

/// isDefinedOutsideFunctionOrMethod - This predicate returns true if this
/// scoped decl is defined outside the current function or method.  This is
/// roughly global variables and functions, but also handles enums (which
/// could be defined inside or outside a function etc).
bool isDefinedOutsideFunctionOrMethod() const {
  return getParentFunctionOrMethod() == nullptr;
}

/// Determine whether a substitution into this declaration would occur as
/// part of a substitution into a dependent local scope. Such a substitution
/// transitively substitutes into all constructs nested within this
/// declaration.
///
/// This recognizes non-defining declarations as well as members of local
/// classes and lambdas:
/// \code
///     template<typename T> void foo() { void bar(); }
///     template<typename T> void foo2() { class ABC { void bar(); }; }
///     template<typename T> inline int x = [](){ return 0; }();
/// \endcode
bool isInLocalScopeForInstantiation() const;

/// If this decl is defined inside a function/method/block it returns
/// the corresponding DeclContext, otherwise it returns null.
const DeclContext *getParentFunctionOrMethod() const;
DeclContext *getParentFunctionOrMethod() {
  return const_cast<DeclContext*>(
                  const_cast<const Decl*>(this)->getParentFunctionOrMethod());
}

/// Retrieves the "canonical" declaration of the given declaration.
virtual Decl *getCanonicalDecl() { return this; }
const Decl *getCanonicalDecl() const {
  return const_cast<Decl*>(this)->getCanonicalDecl();
}

/// Whether this particular Decl is a canonical one.
bool isCanonicalDecl() const { return getCanonicalDecl() == this; }

909protected:
/// Returns the next redeclaration or itself if this is the only decl.
///
/// Decl subclasses that can be redeclared should override this method so that
/// Decl::redecl_iterator can iterate over them.
virtual Decl *getNextRedeclarationImpl() { return this; }

/// Implementation of getPreviousDecl(), to be overridden by any
/// subclass that has a redeclaration chain.
virtual Decl *getPreviousDeclImpl() { return nullptr; }

/// Implementation of getMostRecentDecl(), to be overridden by any
/// subclass that has a redeclaration chain.
virtual Decl *getMostRecentDeclImpl() { return this; }

924public:
/// Iterates through all the redeclarations of the same decl.
class redecl_iterator {
  /// Current - The current declaration.
  Decl *Current = nullptr;
  Decl *Starter;

public:
  using value_type = Decl *;
  using reference = const value_type &;
  using pointer = const value_type *;
  using iterator_category = std::forward_iterator_tag;
  using difference_type = std::ptrdiff_t;

  redecl_iterator() = default;
  explicit redecl_iterator(Decl *C) : Current(C), Starter(C) {}

  reference operator*() const { return Current; }
  value_type operator->() const { return Current; }

  redecl_iterator& operator++() {
    assert(Current && "Advancing while iterator has reached end")((Current && "Advancing while iterator has reached end"
) ? static_cast<void> (0) : __assert_fail ("Current && \"Advancing while iterator has reached end\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 945, __PRETTY_FUNCTION__));
    // Get either previous decl or latest decl.
    Decl *Next = Current->getNextRedeclarationImpl();
    assert(Next && "Should return next redeclaration or itself, never null!")((Next && "Should return next redeclaration or itself, never null!"
) ? static_cast<void> (0) : __assert_fail ("Next && \"Should return next redeclaration or itself, never null!\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 948, __PRETTY_FUNCTION__));
    Current = (Next != Starter) ? Next : nullptr;
    return *this;
  }

  redecl_iterator operator++(int) {
    redecl_iterator tmp(*this);
    ++(*this);
    return tmp;
  }

  friend bool operator==(redecl_iterator x, redecl_iterator y) {
    return x.Current == y.Current;
  }

  friend bool operator!=(redecl_iterator x, redecl_iterator y) {
    return x.Current != y.Current;
  }
};

using redecl_range = llvm::iterator_range<redecl_iterator>;

/// Returns an iterator range for all the redeclarations of the same
/// decl. It will iterate at least once (when this decl is the only one).
redecl_range redecls() const {
  return redecl_range(redecls_begin(), redecls_end());
}

redecl_iterator redecls_begin() const {
  return redecl_iterator(const_cast<Decl *>(this));
}

redecl_iterator redecls_end() const { return redecl_iterator(); }

/// Retrieve the previous declaration that declares the same entity
/// as this declaration, or NULL if there is no previous declaration.
Decl *getPreviousDecl() { return getPreviousDeclImpl(); }

/// Retrieve the previous declaration that declares the same entity
/// as this declaration, or NULL if there is no previous declaration.
const Decl *getPreviousDecl() const {
  return const_cast<Decl *>(this)->getPreviousDeclImpl();
}

/// True if this is the first declaration in its redeclaration chain.
bool isFirstDecl() const {
  return getPreviousDecl() == nullptr;
}

/// Retrieve the most recent declaration that declares the same entity
/// as this declaration (which may be this declaration).
Decl *getMostRecentDecl() { return getMostRecentDeclImpl(); }

/// Retrieve the most recent declaration that declares the same entity
/// as this declaration (which may be this declaration).
const Decl *getMostRecentDecl() const {
  return const_cast<Decl *>(this)->getMostRecentDeclImpl();
}

/// getBody - If this Decl represents a declaration for a body of code,
///  such as a function or method definition, this method returns the
///  top-level Stmt* of that body.  Otherwise this method returns null.
virtual Stmt* getBody() const { return nullptr; }

/// Returns true if this \c Decl represents a declaration for a body of
/// code, such as a function or method definition.
/// Note that \c hasBody can also return true if any redeclaration of this
/// \c Decl represents a declaration for a body of code.
virtual bool hasBody() const { return getBody() != nullptr; }

/// getBodyRBrace - Gets the right brace of the body, if a body exists.
/// This works whether the body is a CompoundStmt or a CXXTryStmt.
SourceLocation getBodyRBrace() const;

// global temp stats (until we have a per-module visitor)
static void add(Kind k);
static void EnableStatistics();
static void PrintStats();

/// isTemplateParameter - Determines whether this declaration is a
/// template parameter.
bool isTemplateParameter() const;

/// isTemplateParameter - Determines whether this declaration is a
/// template parameter pack.
bool isTemplateParameterPack() const;

/// Whether this declaration is a parameter pack.
bool isParameterPack() const;

/// returns true if this declaration is a template
bool isTemplateDecl() const;

/// Whether this declaration is a function or function template.
bool isFunctionOrFunctionTemplate() const {
  return (DeclKind >= Decl::firstFunction &&
          DeclKind <= Decl::lastFunction) ||
         DeclKind == FunctionTemplate;
}

/// If this is a declaration that describes some template, this
/// method returns that template declaration.
///
/// Note that this returns nullptr for partial specializations, because they
/// are not modeled as TemplateDecls. Use getDescribedTemplateParams to handle
/// those cases.
TemplateDecl *getDescribedTemplate() const;

/// If this is a declaration that describes some template or partial
/// specialization, this returns the corresponding template parameter list.
const TemplateParameterList *getDescribedTemplateParams() const;

/// Returns the function itself, or the templated function if this is a
/// function template.
FunctionDecl *getAsFunction() LLVM_READONLY__attribute__((__pure__));

const FunctionDecl *getAsFunction() const {
  return const_cast<Decl *>(this)->getAsFunction();
}

/// Changes the namespace of this declaration to reflect that it's
/// a function-local extern declaration.
///
/// These declarations appear in the lexical context of the extern
/// declaration, but in the semantic context of the enclosing namespace
/// scope.
void setLocalExternDecl() {
  Decl *Prev = getPreviousDecl();
  IdentifierNamespace &= ~IDNS_Ordinary;

  // It's OK for the declaration to still have the "invisible friend" flag or
  // the "conflicts with tag declarations in this scope" flag for the outer
  // scope.
  assert((IdentifierNamespace & ~(IDNS_OrdinaryFriend | IDNS_Tag)) == 0 &&(((IdentifierNamespace & ~(IDNS_OrdinaryFriend | IDNS_Tag
)) == 0 && "namespace is not ordinary") ? static_cast
<void> (0) : __assert_fail ("(IdentifierNamespace & ~(IDNS_OrdinaryFriend | IDNS_Tag)) == 0 && \"namespace is not ordinary\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1082, __PRETTY_FUNCTION__))
         "namespace is not ordinary")(((IdentifierNamespace & ~(IDNS_OrdinaryFriend | IDNS_Tag
)) == 0 && "namespace is not ordinary") ? static_cast
<void> (0) : __assert_fail ("(IdentifierNamespace & ~(IDNS_OrdinaryFriend | IDNS_Tag)) == 0 && \"namespace is not ordinary\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1082, __PRETTY_FUNCTION__));

  IdentifierNamespace |= IDNS_LocalExtern;
  if (Prev && Prev->getIdentifierNamespace() & IDNS_Ordinary)
    IdentifierNamespace |= IDNS_Ordinary;
}

/// Determine whether this is a block-scope declaration with linkage.
/// This will either be a local variable declaration declared 'extern', or a
/// local function declaration.
bool isLocalExternDecl() {
  return IdentifierNamespace & IDNS_LocalExtern;
}

/// Changes the namespace of this declaration to reflect that it's
/// the object of a friend declaration.
///
/// These declarations appear in the lexical context of the friending
/// class, but in the semantic context of the actual entity.  This property
/// applies only to a specific decl object;  other redeclarations of the
/// same entity may not (and probably don't) share this property.
void setObjectOfFriendDecl(bool PerformFriendInjection = false) {
  unsigned OldNS = IdentifierNamespace;
  assert((OldNS & (IDNS_Tag | IDNS_Ordinary |(((OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend
 | IDNS_LocalExtern | IDNS_NonMemberOperator)) && "namespace includes neither ordinary nor tag"
) ? static_cast<void> (0) : __assert_fail ("(OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes neither ordinary nor tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1108, __PRETTY_FUNCTION__))
                   IDNS_TagFriend | IDNS_OrdinaryFriend |(((OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend
 | IDNS_LocalExtern | IDNS_NonMemberOperator)) && "namespace includes neither ordinary nor tag"
) ? static_cast<void> (0) : __assert_fail ("(OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes neither ordinary nor tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1108, __PRETTY_FUNCTION__))
                   IDNS_LocalExtern | IDNS_NonMemberOperator)) &&(((OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend
 | IDNS_LocalExtern | IDNS_NonMemberOperator)) && "namespace includes neither ordinary nor tag"
) ? static_cast<void> (0) : __assert_fail ("(OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes neither ordinary nor tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1108, __PRETTY_FUNCTION__))
         "namespace includes neither ordinary nor tag")(((OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend
 | IDNS_LocalExtern | IDNS_NonMemberOperator)) && "namespace includes neither ordinary nor tag"
) ? static_cast<void> (0) : __assert_fail ("(OldNS & (IDNS_Tag | IDNS_Ordinary | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes neither ordinary nor tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1108, __PRETTY_FUNCTION__));
  assert(!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type |((!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend
 | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator
)) && "namespace includes other than ordinary or tag"
) ? static_cast<void> (0) : __assert_fail ("!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes other than ordinary or tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1112, __PRETTY_FUNCTION__))
                     IDNS_TagFriend | IDNS_OrdinaryFriend |((!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend
 | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator
)) && "namespace includes other than ordinary or tag"
) ? static_cast<void> (0) : __assert_fail ("!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes other than ordinary or tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1112, __PRETTY_FUNCTION__))
                     IDNS_LocalExtern | IDNS_NonMemberOperator)) &&((!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend
 | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator
)) && "namespace includes other than ordinary or tag"
) ? static_cast<void> (0) : __assert_fail ("!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes other than ordinary or tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1112, __PRETTY_FUNCTION__))
         "namespace includes other than ordinary or tag")((!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend
 | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator
)) && "namespace includes other than ordinary or tag"
) ? static_cast<void> (0) : __assert_fail ("!(OldNS & ~(IDNS_Tag | IDNS_Ordinary | IDNS_Type | IDNS_TagFriend | IDNS_OrdinaryFriend | IDNS_LocalExtern | IDNS_NonMemberOperator)) && \"namespace includes other than ordinary or tag\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1112, __PRETTY_FUNCTION__));

  Decl *Prev = getPreviousDecl();
  IdentifierNamespace &= ~(IDNS_Ordinary | IDNS_Tag | IDNS_Type);

  if (OldNS & (IDNS_Tag | IDNS_TagFriend)) {
    IdentifierNamespace |= IDNS_TagFriend;
    if (PerformFriendInjection ||
        (Prev && Prev->getIdentifierNamespace() & IDNS_Tag))
      IdentifierNamespace |= IDNS_Tag | IDNS_Type;
  }

  if (OldNS & (IDNS_Ordinary | IDNS_OrdinaryFriend |
               IDNS_LocalExtern | IDNS_NonMemberOperator)) {
    IdentifierNamespace |= IDNS_OrdinaryFriend;
    if (PerformFriendInjection ||
        (Prev && Prev->getIdentifierNamespace() & IDNS_Ordinary))
      IdentifierNamespace |= IDNS_Ordinary;
  }
}

enum FriendObjectKind {
  FOK_None,      ///< Not a friend object.
  FOK_Declared,  ///< A friend of a previously-declared entity.
  FOK_Undeclared ///< A friend of a previously-undeclared entity.
};

/// Determines whether this declaration is the object of a
/// friend declaration and, if so, what kind.
///
/// There is currently no direct way to find the associated FriendDecl.
FriendObjectKind getFriendObjectKind() const {
  unsigned mask =
      (IdentifierNamespace & (IDNS_TagFriend | IDNS_OrdinaryFriend));
  if (!mask) return FOK_None;
  return (IdentifierNamespace & (IDNS_Tag | IDNS_Ordinary) ? FOK_Declared
                                                           : FOK_Undeclared);
}

/// Specifies that this declaration is a C++ overloaded non-member.
void setNonMemberOperator() {
  assert(getKind() == Function || getKind() == FunctionTemplate)((getKind() == Function || getKind() == FunctionTemplate) ? static_cast
<void> (0) : __assert_fail ("getKind() == Function || getKind() == FunctionTemplate"
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1153, __PRETTY_FUNCTION__));
  assert((IdentifierNamespace & IDNS_Ordinary) &&(((IdentifierNamespace & IDNS_Ordinary) && "visible non-member operators should be in ordinary namespace"
) ? static_cast<void> (0) : __assert_fail ("(IdentifierNamespace & IDNS_Ordinary) && \"visible non-member operators should be in ordinary namespace\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1155, __PRETTY_FUNCTION__))
         "visible non-member operators should be in ordinary namespace")(((IdentifierNamespace & IDNS_Ordinary) && "visible non-member operators should be in ordinary namespace"
) ? static_cast<void> (0) : __assert_fail ("(IdentifierNamespace & IDNS_Ordinary) && \"visible non-member operators should be in ordinary namespace\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1155, __PRETTY_FUNCTION__));
  IdentifierNamespace |= IDNS_NonMemberOperator;
}

static bool classofKind(Kind K) { return true; }
static DeclContext *castToDeclContext(const Decl *);
static Decl *castFromDeclContext(const DeclContext *);

void print(raw_ostream &Out, unsigned Indentation = 0,
           bool PrintInstantiation = false) const;
void print(raw_ostream &Out, const PrintingPolicy &Policy,
           unsigned Indentation = 0, bool PrintInstantiation = false) const;
static void printGroup(Decl** Begin, unsigned NumDecls,
                       raw_ostream &Out, const PrintingPolicy &Policy,
                       unsigned Indentation = 0);

// Debuggers don't usually respect default arguments.
void dump() const;

// Same as dump(), but forces color printing.
void dumpColor() const;

void dump(raw_ostream &Out, bool Deserialize = false,
          ASTDumpOutputFormat OutputFormat = ADOF_Default) const;

/// \return Unique reproducible object identifier
int64_t getID() const;

/// Looks through the Decl's underlying type to extract a FunctionType
/// when possible. Will return null if the type underlying the Decl does not
/// have a FunctionType.
const FunctionType *getFunctionType(bool BlocksToo = true) const;

1188private:
void setAttrsImpl(const AttrVec& Attrs, ASTContext &Ctx);
void setDeclContextsImpl(DeclContext *SemaDC, DeclContext *LexicalDC,
                         ASTContext &Ctx);

1193protected:
ASTMutationListener *getASTMutationListener() const;
1195};

1197/// Determine whether two declarations declare the same entity.
1198inline bool declaresSameEntity(const Decl *D1, const Decl *D2) {
if (!D1 || !D2)
  return false;

if (D1 == D2)
  return true;

return D1->getCanonicalDecl() == D2->getCanonicalDecl();
1206}

1208/// PrettyStackTraceDecl - If a crash occurs, indicate that it happened when
1209/// doing something to a specific decl.
1210class PrettyStackTraceDecl : public llvm::PrettyStackTraceEntry {
const Decl *TheDecl;
SourceLocation Loc;
SourceManager &SM;
const char *Message;

1216public:
PrettyStackTraceDecl(const Decl *theDecl, SourceLocation L,
                     SourceManager &sm, const char *Msg)
    : TheDecl(theDecl), Loc(L), SM(sm), Message(Msg) {}

void print(raw_ostream &OS) const override;
1222};
1223} // namespace clang

1225// Required to determine the layout of the PointerUnion<NamedDecl*> before
1226// seeing the NamedDecl definition being first used in DeclListNode::operator*.
1227namespace llvm {
template <> struct PointerLikeTypeTraits<::clang::NamedDecl *> {
  static inline void *getAsVoidPointer(::clang::NamedDecl *P) { return P; }
  static inline ::clang::NamedDecl *getFromVoidPointer(void *P) {
    return static_cast<::clang::NamedDecl *>(P);
  }
  static constexpr int NumLowBitsAvailable = 3;
};
1235}

1237namespace clang {
1238/// A list storing NamedDecls in the lookup tables.
1239class DeclListNode {
friend class ASTContext; // allocate, deallocate nodes.
friend class StoredDeclsList;
1242public:
using Decls = llvm::PointerUnion<NamedDecl*, DeclListNode*>;
class iterator {
  friend class DeclContextLookupResult;
  friend class StoredDeclsList;

  Decls Ptr;
  iterator(Decls Node) : Ptr(Node) { }
public:
  using difference_type = ptrdiff_t;
  using value_type = NamedDecl*;
  using pointer = void;
  using reference = value_type;
  using iterator_category = std::forward_iterator_tag;

  iterator() = default;

  reference operator*() const {
    assert(Ptr && "dereferencing end() iterator")((Ptr && "dereferencing end() iterator") ? static_cast
<void> (0) : __assert_fail ("Ptr && \"dereferencing end() iterator\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1260, __PRETTY_FUNCTION__));
    if (DeclListNode *CurNode = Ptr.dyn_cast<DeclListNode*>())
      return CurNode->D;
    return Ptr.get<NamedDecl*>();
  }
  void operator->() const { } // Unsupported.
  bool operator==(const iterator &X) const { return Ptr == X.Ptr; }
  bool operator!=(const iterator &X) const { return Ptr != X.Ptr; }
  inline iterator &operator++() { // ++It
    assert(!Ptr.isNull() && "Advancing empty iterator")((!Ptr.isNull() && "Advancing empty iterator") ? static_cast
<void> (0) : __assert_fail ("!Ptr.isNull() && \"Advancing empty iterator\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 1269, __PRETTY_FUNCTION__));

    if (DeclListNode *CurNode = Ptr.dyn_cast<DeclListNode*>())
      Ptr = CurNode->Rest;
    else
      Ptr = nullptr;
    return *this;
  }
  iterator operator++(int) { // It++
    iterator temp = *this;
    ++(*this);
    return temp;
  }
  // Enables the pattern for (iterator I =..., E = I.end(); I != E; ++I)
  iterator end() { return iterator(); }
};
1285private:
NamedDecl *D = nullptr;
Decls Rest = nullptr;
DeclListNode(NamedDecl *ND) : D(ND) {}
1289};

1291/// The results of name lookup within a DeclContext.
1292class DeclContextLookupResult {
using Decls = DeclListNode::Decls;

/// When in collection form, this is what the Data pointer points to.
Decls Result;

1298public:
DeclContextLookupResult() = default;
DeclContextLookupResult(Decls Result) : Result(Result) {}

using iterator = DeclListNode::iterator;
using const_iterator = iterator;
using reference = iterator::reference;

iterator begin() { return iterator(Result); }
iterator end() { return iterator(); }
const_iterator begin() const {
  return const_cast<DeclContextLookupResult*>(this)->begin();
}
const_iterator end() const { return iterator(); }

bool empty() const { return Result.isNull();  }
bool isSingleResult() const { return Result.dyn_cast<NamedDecl*>(); }
reference front() const { return *begin(); }

// Find the first declaration of the given type in the list. Note that this
// is not in general the earliest-declared declaration, and should only be
// used when it's not possible for there to be more than one match or where
// it doesn't matter which one is found.
template<class T> T *find_first() const {
  for (auto *D : *this)
    if (T *Decl = dyn_cast<T>(D))
      return Decl;

  return nullptr;
}
1328};

1330/// DeclContext - This is used only as base class of specific decl types that
1331/// can act as declaration contexts. These decls are (only the top classes
1332/// that directly derive from DeclContext are mentioned, not their subclasses):
1333///
1334///   TranslationUnitDecl
1335///   ExternCContext
1336///   NamespaceDecl
1337///   TagDecl
1338///   OMPDeclareReductionDecl
1339///   OMPDeclareMapperDecl
1340///   FunctionDecl
1341///   ObjCMethodDecl
1342///   ObjCContainerDecl
1343///   LinkageSpecDecl
1344///   ExportDecl
1345///   BlockDecl
1346///   CapturedDecl
1347class DeclContext {
/// For makeDeclVisibleInContextImpl
friend class ASTDeclReader;
/// For reconcileExternalVisibleStorage, CreateStoredDeclsMap,
/// hasNeedToReconcileExternalVisibleStorage
friend class ExternalASTSource;
/// For CreateStoredDeclsMap
friend class DependentDiagnostic;
/// For hasNeedToReconcileExternalVisibleStorage,
/// hasLazyLocalLexicalLookups, hasLazyExternalLexicalLookups
friend class ASTWriter;

// We use uint64_t in the bit-fields below since some bit-fields
// cross the unsigned boundary and this breaks the packing.

/// Stores the bits used by DeclContext.
/// If modified NumDeclContextBit, the ctor of DeclContext and the accessor
/// methods in DeclContext should be updated appropriately.
class DeclContextBitfields {
  friend class DeclContext;
  /// DeclKind - This indicates which class this is.
  uint64_t DeclKind : 7;

  /// Whether this declaration context also has some external
  /// storage that contains additional declarations that are lexically
  /// part of this context.
  mutable uint64_t ExternalLexicalStorage : 1;

  /// Whether this declaration context also has some external
  /// storage that contains additional declarations that are visible
  /// in this context.
  mutable uint64_t ExternalVisibleStorage : 1;

  /// Whether this declaration context has had externally visible
  /// storage added since the last lookup. In this case, \c LookupPtr's
  /// invariant may not hold and needs to be fixed before we perform
  /// another lookup.
  mutable uint64_t NeedToReconcileExternalVisibleStorage : 1;

  /// If \c true, this context may have local lexical declarations
  /// that are missing from the lookup table.
  mutable uint64_t HasLazyLocalLexicalLookups : 1;

  /// If \c true, the external source may have lexical declarations
  /// that are missing from the lookup table.
  mutable uint64_t HasLazyExternalLexicalLookups : 1;

  /// If \c true, lookups should only return identifier from
  /// DeclContext scope (for example TranslationUnit). Used in
  /// LookupQualifiedName()
  mutable uint64_t UseQualifiedLookup : 1;
};

/// Number of bits in DeclContextBitfields.
enum { NumDeclContextBits = 13 };

/// Stores the bits used by TagDecl.
/// If modified NumTagDeclBits and the accessor
/// methods in TagDecl should be updated appropriately.
class TagDeclBitfields {
  friend class TagDecl;
  /// For the bits in DeclContextBitfields
  uint64_t : NumDeclContextBits;

  /// The TagKind enum.
  uint64_t TagDeclKind : 3;

  /// True if this is a definition ("struct foo {};"), false if it is a
  /// declaration ("struct foo;").  It is not considered a definition
  /// until the definition has been fully processed.
  uint64_t IsCompleteDefinition : 1;

  /// True if this is currently being defined.
  uint64_t IsBeingDefined : 1;

  /// True if this tag declaration is "embedded" (i.e., defined or declared
  /// for the very first time) in the syntax of a declarator.
  uint64_t IsEmbeddedInDeclarator : 1;

  /// True if this tag is free standing, e.g. "struct foo;".
  uint64_t IsFreeStanding : 1;

  /// Indicates whether it is possible for declarations of this kind
  /// to have an out-of-date definition.
  ///
  /// This option is only enabled when modules are enabled.
  uint64_t MayHaveOutOfDateDef : 1;

  /// Has the full definition of this type been required by a use somewhere in
  /// the TU.
  uint64_t IsCompleteDefinitionRequired : 1;
};

/// Number of non-inherited bits in TagDeclBitfields.
enum { NumTagDeclBits = 9 };

/// Stores the bits used by EnumDecl.
/// If modified NumEnumDeclBit and the accessor
/// methods in EnumDecl should be updated appropriately.
class EnumDeclBitfields {
  friend class EnumDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;
  /// For the bits in TagDeclBitfields.
  uint64_t : NumTagDeclBits;

  /// Width in bits required to store all the non-negative
  /// enumerators of this enum.
  uint64_t NumPositiveBits : 8;

  /// Width in bits required to store all the negative
  /// enumerators of this enum.
  uint64_t NumNegativeBits : 8;

  /// True if this tag declaration is a scoped enumeration. Only
  /// possible in C++11 mode.
  uint64_t IsScoped : 1;

  /// If this tag declaration is a scoped enum,
  /// then this is true if the scoped enum was declared using the class
  /// tag, false if it was declared with the struct tag. No meaning is
  /// associated if this tag declaration is not a scoped enum.
  uint64_t IsScopedUsingClassTag : 1;

  /// True if this is an enumeration with fixed underlying type. Only
  /// possible in C++11, Microsoft extensions, or Objective C mode.
  uint64_t IsFixed : 1;

  /// True if a valid hash is stored in ODRHash.
  uint64_t HasODRHash : 1;
};

/// Number of non-inherited bits in EnumDeclBitfields.
enum { NumEnumDeclBits = 20 };

/// Stores the bits used by RecordDecl.
/// If modified NumRecordDeclBits and the accessor
/// methods in RecordDecl should be updated appropriately.
class RecordDeclBitfields {
  friend class RecordDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;
  /// For the bits in TagDeclBitfields.
  uint64_t : NumTagDeclBits;

  /// This is true if this struct ends with a flexible
  /// array member (e.g. int X[]) or if this union contains a struct that does.
  /// If so, this cannot be contained in arrays or other structs as a member.
  uint64_t HasFlexibleArrayMember : 1;

  /// Whether this is the type of an anonymous struct or union.
  uint64_t AnonymousStructOrUnion : 1;

  /// This is true if this struct has at least one member
  /// containing an Objective-C object pointer type.
  uint64_t HasObjectMember : 1;

  /// This is true if struct has at least one member of
  /// 'volatile' type.
  uint64_t HasVolatileMember : 1;

  /// Whether the field declarations of this record have been loaded
  /// from external storage. To avoid unnecessary deserialization of
  /// methods/nested types we allow deserialization of just the fields
  /// when needed.
  mutable uint64_t LoadedFieldsFromExternalStorage : 1;

  /// Basic properties of non-trivial C structs.
  uint64_t NonTrivialToPrimitiveDefaultInitialize : 1;
  uint64_t NonTrivialToPrimitiveCopy : 1;
  uint64_t NonTrivialToPrimitiveDestroy : 1;

  /// The following bits indicate whether this is or contains a C union that
  /// is non-trivial to default-initialize, destruct, or copy. These bits
  /// imply the associated basic non-triviality predicates declared above.
  uint64_t HasNonTrivialToPrimitiveDefaultInitializeCUnion : 1;
  uint64_t HasNonTrivialToPrimitiveDestructCUnion : 1;
  uint64_t HasNonTrivialToPrimitiveCopyCUnion : 1;

  /// Indicates whether this struct is destroyed in the callee.
  uint64_t ParamDestroyedInCallee : 1;

  /// Represents the way this type is passed to a function.
  uint64_t ArgPassingRestrictions : 2;
};

/// Number of non-inherited bits in RecordDeclBitfields.
enum { NumRecordDeclBits = 14 };

/// Stores the bits used by OMPDeclareReductionDecl.
/// If modified NumOMPDeclareReductionDeclBits and the accessor
/// methods in OMPDeclareReductionDecl should be updated appropriately.
class OMPDeclareReductionDeclBitfields {
  friend class OMPDeclareReductionDecl;
  /// For the bits in DeclContextBitfields
  uint64_t : NumDeclContextBits;

  /// Kind of initializer,
  /// function call or omp_priv<init_expr> initializtion.
  uint64_t InitializerKind : 2;
};

/// Number of non-inherited bits in OMPDeclareReductionDeclBitfields.
enum { NumOMPDeclareReductionDeclBits = 2 };

/// Stores the bits used by FunctionDecl.
/// If modified NumFunctionDeclBits and the accessor
/// methods in FunctionDecl and CXXDeductionGuideDecl
/// (for IsCopyDeductionCandidate) should be updated appropriately.
class FunctionDeclBitfields {
  friend class FunctionDecl;
  /// For IsCopyDeductionCandidate
  friend class CXXDeductionGuideDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;

  uint64_t SClass : 3;
  uint64_t IsInline : 1;
  uint64_t IsInlineSpecified : 1;

  uint64_t IsVirtualAsWritten : 1;
  uint64_t IsPure : 1;
  uint64_t HasInheritedPrototype : 1;
  uint64_t HasWrittenPrototype : 1;
  uint64_t IsDeleted : 1;
  /// Used by CXXMethodDecl
  uint64_t IsTrivial : 1;

  /// This flag indicates whether this function is trivial for the purpose of
  /// calls. This is meaningful only when this function is a copy/move
  /// constructor or a destructor.
  uint64_t IsTrivialForCall : 1;

  uint64_t IsDefaulted : 1;
  uint64_t IsExplicitlyDefaulted : 1;
  uint64_t HasDefaultedFunctionInfo : 1;
  uint64_t HasImplicitReturnZero : 1;
  uint64_t IsLateTemplateParsed : 1;

  /// Kind of contexpr specifier as defined by ConstexprSpecKind.
  uint64_t ConstexprKind : 2;
  uint64_t InstantiationIsPending : 1;

  /// Indicates if the function uses __try.
  uint64_t UsesSEHTry : 1;

  /// Indicates if the function was a definition
  /// but its body was skipped.
  uint64_t HasSkippedBody : 1;

  /// Indicates if the function declaration will
  /// have a body, once we're done parsing it.
  uint64_t WillHaveBody : 1;

  /// Indicates that this function is a multiversioned
  /// function using attribute 'target'.
  uint64_t IsMultiVersion : 1;

  /// [C++17] Only used by CXXDeductionGuideDecl. Indicates that
  /// the Deduction Guide is the implicitly generated 'copy
  /// deduction candidate' (is used during overload resolution).
  uint64_t IsCopyDeductionCandidate : 1;

  /// Store the ODRHash after first calculation.
  uint64_t HasODRHash : 1;

  /// Indicates if the function uses Floating Point Constrained Intrinsics
  uint64_t UsesFPIntrin : 1;
};

/// Number of non-inherited bits in FunctionDeclBitfields.
enum { NumFunctionDeclBits = 27 };

/// Stores the bits used by CXXConstructorDecl. If modified
/// NumCXXConstructorDeclBits and the accessor
/// methods in CXXConstructorDecl should be updated appropriately.
class CXXConstructorDeclBitfields {
  friend class CXXConstructorDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;
  /// For the bits in FunctionDeclBitfields.
  uint64_t : NumFunctionDeclBits;

  /// 24 bits to fit in the remaining available space.
  /// Note that this makes CXXConstructorDeclBitfields take
  /// exactly 64 bits and thus the width of NumCtorInitializers
  /// will need to be shrunk if some bit is added to NumDeclContextBitfields,
  /// NumFunctionDeclBitfields or CXXConstructorDeclBitfields.
  uint64_t NumCtorInitializers : 21;
  uint64_t IsInheritingConstructor : 1;

  /// Whether this constructor has a trail-allocated explicit specifier.
  uint64_t HasTrailingExplicitSpecifier : 1;
  /// If this constructor does't have a trail-allocated explicit specifier.
  /// Whether this constructor is explicit specified.
  uint64_t IsSimpleExplicit : 1;
};

/// Number of non-inherited bits in CXXConstructorDeclBitfields.
enum {
  NumCXXConstructorDeclBits = 64 - NumDeclContextBits - NumFunctionDeclBits
};

/// Stores the bits used by ObjCMethodDecl.
/// If modified NumObjCMethodDeclBits and the accessor
/// methods in ObjCMethodDecl should be updated appropriately.
class ObjCMethodDeclBitfields {
  friend class ObjCMethodDecl;

  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;

  /// The conventional meaning of this method; an ObjCMethodFamily.
  /// This is not serialized; instead, it is computed on demand and
  /// cached.
  mutable uint64_t Family : ObjCMethodFamilyBitWidth;

  /// instance (true) or class (false) method.
  uint64_t IsInstance : 1;
  uint64_t IsVariadic : 1;

  /// True if this method is the getter or setter for an explicit property.
  uint64_t IsPropertyAccessor : 1;

  /// True if this method is a synthesized property accessor stub.
  uint64_t IsSynthesizedAccessorStub : 1;

  /// Method has a definition.
  uint64_t IsDefined : 1;

  /// Method redeclaration in the same interface.
  uint64_t IsRedeclaration : 1;

  /// Is redeclared in the same interface.
  mutable uint64_t HasRedeclaration : 1;

  /// \@required/\@optional
  uint64_t DeclImplementation : 2;

  /// in, inout, etc.
  uint64_t objcDeclQualifier : 7;

  /// Indicates whether this method has a related result type.
  uint64_t RelatedResultType : 1;

  /// Whether the locations of the selector identifiers are in a
  /// "standard" position, a enum SelectorLocationsKind.
  uint64_t SelLocsKind : 2;

  /// Whether this method overrides any other in the class hierarchy.
  ///
  /// A method is said to override any method in the class's
  /// base classes, its protocols, or its categories' protocols, that has
  /// the same selector and is of the same kind (class or instance).
  /// A method in an implementation is not considered as overriding the same
  /// method in the interface or its categories.
  uint64_t IsOverriding : 1;

  /// Indicates if the method was a definition but its body was skipped.
  uint64_t HasSkippedBody : 1;
};

/// Number of non-inherited bits in ObjCMethodDeclBitfields.
enum { NumObjCMethodDeclBits = 24 };

/// Stores the bits used by ObjCContainerDecl.
/// If modified NumObjCContainerDeclBits and the accessor
/// methods in ObjCContainerDecl should be updated appropriately.
class ObjCContainerDeclBitfields {
  friend class ObjCContainerDecl;
  /// For the bits in DeclContextBitfields
  uint32_t : NumDeclContextBits;

  // Not a bitfield but this saves space.
  // Note that ObjCContainerDeclBitfields is full.
  SourceLocation AtStart;
};

/// Number of non-inherited bits in ObjCContainerDeclBitfields.
/// Note that here we rely on the fact that SourceLocation is 32 bits
/// wide. We check this with the static_assert in the ctor of DeclContext.
enum { NumObjCContainerDeclBits = 64 - NumDeclContextBits };

/// Stores the bits used by LinkageSpecDecl.
/// If modified NumLinkageSpecDeclBits and the accessor
/// methods in LinkageSpecDecl should be updated appropriately.
class LinkageSpecDeclBitfields {
  friend class LinkageSpecDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;

  /// The language for this linkage specification with values
  /// in the enum LinkageSpecDecl::LanguageIDs.
  uint64_t Language : 3;

  /// True if this linkage spec has braces.
  /// This is needed so that hasBraces() returns the correct result while the
  /// linkage spec body is being parsed.  Once RBraceLoc has been set this is
  /// not used, so it doesn't need to be serialized.
  uint64_t HasBraces : 1;
};

/// Number of non-inherited bits in LinkageSpecDeclBitfields.
enum { NumLinkageSpecDeclBits = 4 };

/// Stores the bits used by BlockDecl.
/// If modified NumBlockDeclBits and the accessor
/// methods in BlockDecl should be updated appropriately.
class BlockDeclBitfields {
  friend class BlockDecl;
  /// For the bits in DeclContextBitfields.
  uint64_t : NumDeclContextBits;

  uint64_t IsVariadic : 1;
  uint64_t CapturesCXXThis : 1;
  uint64_t BlockMissingReturnType : 1;
  uint64_t IsConversionFromLambda : 1;

  /// A bit that indicates this block is passed directly to a function as a
  /// non-escaping parameter.
  uint64_t DoesNotEscape : 1;

  /// A bit that indicates whether it's possible to avoid coying this block to
  /// the heap when it initializes or is assigned to a local variable with
  /// automatic storage.
  uint64_t CanAvoidCopyToHeap : 1;
};

/// Number of non-inherited bits in BlockDeclBitfields.
enum { NumBlockDeclBits = 5 };

/// Pointer to the data structure used to lookup declarations
/// within this context (or a DependentStoredDeclsMap if this is a
/// dependent context). We maintain the invariant that, if the map
/// contains an entry for a DeclarationName (and we haven't lazily
/// omitted anything), then it contains all relevant entries for that
/// name (modulo the hasExternalDecls() flag).
mutable StoredDeclsMap *LookupPtr = nullptr;

1786protected:
/// This anonymous union stores the bits belonging to DeclContext and classes
/// deriving from it. The goal is to use otherwise wasted
/// space in DeclContext to store data belonging to derived classes.
/// The space saved is especially significient when pointers are aligned
/// to 8 bytes. In this case due to alignment requirements we have a
/// little less than 8 bytes free in DeclContext which we can use.
/// We check that none of the classes in this union is larger than
/// 8 bytes with static_asserts in the ctor of DeclContext.
union {
  DeclContextBitfields DeclContextBits;
  TagDeclBitfields TagDeclBits;
  EnumDeclBitfields EnumDeclBits;
  RecordDeclBitfields RecordDeclBits;
  OMPDeclareReductionDeclBitfields OMPDeclareReductionDeclBits;
  FunctionDeclBitfields FunctionDeclBits;
  CXXConstructorDeclBitfields CXXConstructorDeclBits;
  ObjCMethodDeclBitfields ObjCMethodDeclBits;
  ObjCContainerDeclBitfields ObjCContainerDeclBits;
  LinkageSpecDeclBitfields LinkageSpecDeclBits;
  BlockDeclBitfields BlockDeclBits;

  static_assert(sizeof(DeclContextBitfields) <= 8,
                "DeclContextBitfields is larger than 8 bytes!");
  static_assert(sizeof(TagDeclBitfields) <= 8,
                "TagDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(EnumDeclBitfields) <= 8,
                "EnumDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(RecordDeclBitfields) <= 8,
                "RecordDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(OMPDeclareReductionDeclBitfields) <= 8,
                "OMPDeclareReductionDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(FunctionDeclBitfields) <= 8,
                "FunctionDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(CXXConstructorDeclBitfields) <= 8,
                "CXXConstructorDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(ObjCMethodDeclBitfields) <= 8,
                "ObjCMethodDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(ObjCContainerDeclBitfields) <= 8,
                "ObjCContainerDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(LinkageSpecDeclBitfields) <= 8,
                "LinkageSpecDeclBitfields is larger than 8 bytes!");
  static_assert(sizeof(BlockDeclBitfields) <= 8,
                "BlockDeclBitfields is larger than 8 bytes!");
};

/// FirstDecl - The first declaration stored within this declaration
/// context.
mutable Decl *FirstDecl = nullptr;

/// LastDecl - The last declaration stored within this declaration
/// context. FIXME: We could probably cache this value somewhere
/// outside of the DeclContext, to reduce the size of DeclContext by
/// another pointer.
mutable Decl *LastDecl = nullptr;

/// Build up a chain of declarations.
///
/// \returns the first/last pair of declarations.
static std::pair<Decl *, Decl *>
BuildDeclChain(ArrayRef<Decl*> Decls, bool FieldsAlreadyLoaded);

DeclContext(Decl::Kind K);

1850public:
~DeclContext();

Decl::Kind getDeclKind() const {
  return static_cast<Decl::Kind>(DeclContextBits.DeclKind);
}

const char *getDeclKindName() const;

/// getParent - Returns the containing DeclContext.
DeclContext *getParent() {
  return cast<Decl>(this)->getDeclContext();
}
const DeclContext *getParent() const {
  return const_cast<DeclContext*>(this)->getParent();
}

/// getLexicalParent - Returns the containing lexical DeclContext. May be
/// different from getParent, e.g.:
///
///   namespace A {
///      struct S;
///   }
///   struct A::S {}; // getParent() == namespace 'A'
///                   // getLexicalParent() == translation unit
///
DeclContext *getLexicalParent() {
  return cast<Decl>(this)->getLexicalDeclContext();
}
const DeclContext *getLexicalParent() const {
  return const_cast<DeclContext*>(this)->getLexicalParent();
}

DeclContext *getLookupParent();

const DeclContext *getLookupParent() const {
  return const_cast<DeclContext*>(this)->getLookupParent();
}

ASTContext &getParentASTContext() const {
  return cast<Decl>(this)->getASTContext();
}

bool isClosure() const { return getDeclKind() == Decl::Block; }

/// Return this DeclContext if it is a BlockDecl. Otherwise, return the
/// innermost enclosing BlockDecl or null if there are no enclosing blocks.
const BlockDecl *getInnermostBlockDecl() const;

bool isObjCContainer() const {
  switch (getDeclKind()) {
  case Decl::ObjCCategory:
  case Decl::ObjCCategoryImpl:
  case Decl::ObjCImplementation:
  case Decl::ObjCInterface:
  case Decl::ObjCProtocol:
    return true;
  default:
    return false;
  }
}

bool isFunctionOrMethod() const {
  switch (getDeclKind()) {
  case Decl::Block:
  case Decl::Captured:
  case Decl::ObjCMethod:
    return true;
  default:
    return getDeclKind() >= Decl::firstFunction &&
           getDeclKind() <= Decl::lastFunction;
  }
}

/// Test whether the context supports looking up names.
bool isLookupContext() const {
  return !isFunctionOrMethod() && getDeclKind() != Decl::LinkageSpec &&
         getDeclKind() != Decl::Export;
}

bool isFileContext() const {
  return getDeclKind() == Decl::TranslationUnit ||
         getDeclKind() == Decl::Namespace;
}

bool isTranslationUnit() const {
  return getDeclKind() == Decl::TranslationUnit;
}

bool isRecord() const {
  return getDeclKind() >= Decl::firstRecord &&
         getDeclKind() <= Decl::lastRecord;
}

bool isNamespace() const { return getDeclKind() == Decl::Namespace; }

bool isStdNamespace() const;

bool isInlineNamespace() const;

/// Determines whether this context is dependent on a
/// template parameter.
bool isDependentContext() const;

/// isTransparentContext - Determines whether this context is a
/// "transparent" context, meaning that the members declared in this
/// context are semantically declared in the nearest enclosing
/// non-transparent (opaque) context but are lexically declared in
/// this context. For example, consider the enumerators of an
/// enumeration type:
/// @code
/// enum E {
///   Val1
/// };
/// @endcode
/// Here, E is a transparent context, so its enumerator (Val1) will
/// appear (semantically) that it is in the same context of E.
/// Examples of transparent contexts include: enumerations (except for
/// C++0x scoped enums), and C++ linkage specifications.
bool isTransparentContext() const;

/// Determines whether this context or some of its ancestors is a
/// linkage specification context that specifies C linkage.
bool isExternCContext() const;

/// Retrieve the nearest enclosing C linkage specification context.
const LinkageSpecDecl *getExternCContext() const;

/// Determines whether this context or some of its ancestors is a
/// linkage specification context that specifies C++ linkage.
bool isExternCXXContext() const;

/// Determine whether this declaration context is equivalent
/// to the declaration context DC.
bool Equals(const DeclContext *DC) const {
  return DC && this->getPrimaryContext() == DC->getPrimaryContext();
}

/// Determine whether this declaration context encloses the
/// declaration context DC.
bool Encloses(const DeclContext *DC) const;

/// Find the nearest non-closure ancestor of this context,
/// i.e. the innermost semantic parent of this context which is not
/// a closure.  A context may be its own non-closure ancestor.
Decl *getNonClosureAncestor();
const Decl *getNonClosureAncestor() const {
  return const_cast<DeclContext*>(this)->getNonClosureAncestor();
}

/// getPrimaryContext - There may be many different
/// declarations of the same entity (including forward declarations
/// of classes, multiple definitions of namespaces, etc.), each with
/// a different set of declarations. This routine returns the
/// "primary" DeclContext structure, which will contain the
/// information needed to perform name lookup into this context.
DeclContext *getPrimaryContext();
const DeclContext *getPrimaryContext() const {
  return const_cast<DeclContext*>(this)->getPrimaryContext();
}

/// getRedeclContext - Retrieve the context in which an entity conflicts with
/// other entities of the same name, or where it is a redeclaration if the
/// two entities are compatible. This skips through transparent contexts.
DeclContext *getRedeclContext();
const DeclContext *getRedeclContext() const {
  return const_cast<DeclContext *>(this)->getRedeclContext();
}

/// Retrieve the nearest enclosing namespace context.
DeclContext *getEnclosingNamespaceContext();
const DeclContext *getEnclosingNamespaceContext() const {
  return const_cast<DeclContext *>(this)->getEnclosingNamespaceContext();
}

/// Retrieve the outermost lexically enclosing record context.
RecordDecl *getOuterLexicalRecordContext();
const RecordDecl *getOuterLexicalRecordContext() const {
  return const_cast<DeclContext *>(this)->getOuterLexicalRecordContext();
}

/// Test if this context is part of the enclosing namespace set of
/// the context NS, as defined in C++0x [namespace.def]p9. If either context
/// isn't a namespace, this is equivalent to Equals().
///
/// The enclosing namespace set of a namespace is the namespace and, if it is
/// inline, its enclosing namespace, recursively.
bool InEnclosingNamespaceSetOf(const DeclContext *NS) const;

/// Collects all of the declaration contexts that are semantically
/// connected to this declaration context.
///
/// For declaration contexts that have multiple semantically connected but
/// syntactically distinct contexts, such as C++ namespaces, this routine
/// retrieves the complete set of such declaration contexts in source order.
/// For example, given:
///
/// \code
/// namespace N {
///   int x;
/// }
/// namespace N {
///   int y;
/// }
/// \endcode
///
/// The \c Contexts parameter will contain both definitions of N.
///
/// \param Contexts Will be cleared and set to the set of declaration
/// contexts that are semanticaly connected to this declaration context,
/// in source order, including this context (which may be the only result,
/// for non-namespace contexts).
void collectAllContexts(SmallVectorImpl<DeclContext *> &Contexts);

/// decl_iterator - Iterates through the declarations stored
/// within this context.
class decl_iterator {
  /// Current - The current declaration.
  Decl *Current = nullptr;

public:
  using value_type = Decl *;
  using reference = const value_type &;
  using pointer = const value_type *;
  using iterator_category = std::forward_iterator_tag;
  using difference_type = std::ptrdiff_t;

  decl_iterator() = default;
  explicit decl_iterator(Decl *C) : Current(C) {}

  reference operator*() const { return Current; }

  // This doesn't meet the iterator requirements, but it's convenient
  value_type operator->() const { return Current; }

  decl_iterator& operator++() {
    Current = Current->getNextDeclInContext();
    return *this;
  }

  decl_iterator operator++(int) {
    decl_iterator tmp(*this);
    ++(*this);
    return tmp;
  }

  friend bool operator==(decl_iterator x, decl_iterator y) {
    return x.Current == y.Current;
  }

  friend bool operator!=(decl_iterator x, decl_iterator y) {
    return x.Current != y.Current;
  }
};

using decl_range = llvm::iterator_range<decl_iterator>;

/// decls_begin/decls_end - Iterate over the declarations stored in
/// this context.
decl_range decls() const { return decl_range(decls_begin(), decls_end()); }
decl_iterator decls_begin() const;
decl_iterator decls_end() const { return decl_iterator(); }
bool decls_empty() const;

/// noload_decls_begin/end - Iterate over the declarations stored in this
/// context that are currently loaded; don't attempt to retrieve anything
/// from an external source.
decl_range noload_decls() const {
  return decl_range(noload_decls_begin(), noload_decls_end());
}
decl_iterator noload_decls_begin() const { return decl_iterator(FirstDecl); }
decl_iterator noload_decls_end() const { return decl_iterator(); }

/// specific_decl_iterator - Iterates over a subrange of
/// declarations stored in a DeclContext, providing only those that
/// are of type SpecificDecl (or a class derived from it). This
/// iterator is used, for example, to provide iteration over just
/// the fields within a RecordDecl (with SpecificDecl = FieldDecl).
template<typename SpecificDecl>
class specific_decl_iterator {
  /// Current - The current, underlying declaration iterator, which
  /// will either be NULL or will point to a declaration of
  /// type SpecificDecl.
  DeclContext::decl_iterator Current;

  /// SkipToNextDecl - Advances the current position up to the next
  /// declaration of type SpecificDecl that also meets the criteria
  /// required by Acceptable.
  void SkipToNextDecl() {
    while (*Current && !isa<SpecificDecl>(*Current))
      ++Current;
  }

public:
  using value_type = SpecificDecl *;
  // TODO: Add reference and pointer types (with some appropriate proxy type)
  // if we ever have a need for them.
  using reference = void;
  using pointer = void;
  using difference_type =
      std::iterator_traits<DeclContext::decl_iterator>::difference_type;
  using iterator_category = std::forward_iterator_tag;

  specific_decl_iterator() = default;

  /// specific_decl_iterator - Construct a new iterator over a
  /// subset of the declarations the range [C,
  /// end-of-declarations). If A is non-NULL, it is a pointer to a
  /// member function of SpecificDecl that should return true for
  /// all of the SpecificDecl instances that will be in the subset
  /// of iterators. For example, if you want Objective-C instance
  /// methods, SpecificDecl will be ObjCMethodDecl and A will be
  /// &ObjCMethodDecl::isInstanceMethod.
  explicit specific_decl_iterator(DeclContext::decl_iterator C) : Current(C) {
    SkipToNextDecl();
  }

  value_type operator*() const { return cast<SpecificDecl>(*Current); }

  // This doesn't meet the iterator requirements, but it's convenient
  value_type operator->() const { return **this; }

  specific_decl_iterator& operator++() {
    ++Current;
    SkipToNextDecl();
    return *this;
  }

  specific_decl_iterator operator++(int) {
    specific_decl_iterator tmp(*this);
    ++(*this);
    return tmp;
  }

  friend bool operator==(const specific_decl_iterator& x,
                         const specific_decl_iterator& y) {
    return x.Current == y.Current;
  }

  friend bool operator!=(const specific_decl_iterator& x,
                         const specific_decl_iterator& y) {
    return x.Current != y.Current;
  }
};

/// Iterates over a filtered subrange of declarations stored
/// in a DeclContext.
///
/// This iterator visits only those declarations that are of type
/// SpecificDecl (or a class derived from it) and that meet some
/// additional run-time criteria. This iterator is used, for
/// example, to provide access to the instance methods within an
/// Objective-C interface (with SpecificDecl = ObjCMethodDecl and
/// Acceptable = ObjCMethodDecl::isInstanceMethod).
template<typename SpecificDecl, bool (SpecificDecl::*Acceptable)() const>
class filtered_decl_iterator {
  /// Current - The current, underlying declaration iterator, which
  /// will either be NULL or will point to a declaration of
  /// type SpecificDecl.
  DeclContext::decl_iterator Current;

  /// SkipToNextDecl - Advances the current position up to the next
  /// declaration of type SpecificDecl that also meets the criteria
  /// required by Acceptable.
  void SkipToNextDecl() {
    while (*Current &&
           (!isa<SpecificDecl>(*Current) ||
            (Acceptable && !(cast<SpecificDecl>(*Current)->*Acceptable)())))
      ++Current;
  }

public:
  using value_type = SpecificDecl *;
  // TODO: Add reference and pointer types (with some appropriate proxy type)
  // if we ever have a need for them.
  using reference = void;
  using pointer = void;
  using difference_type =
      std::iterator_traits<DeclContext::decl_iterator>::difference_type;
  using iterator_category = std::forward_iterator_tag;

  filtered_decl_iterator() = default;

  /// filtered_decl_iterator - Construct a new iterator over a
  /// subset of the declarations the range [C,
  /// end-of-declarations). If A is non-NULL, it is a pointer to a
  /// member function of SpecificDecl that should return true for
  /// all of the SpecificDecl instances that will be in the subset
  /// of iterators. For example, if you want Objective-C instance
  /// methods, SpecificDecl will be ObjCMethodDecl and A will be
  /// &ObjCMethodDecl::isInstanceMethod.
  explicit filtered_decl_iterator(DeclContext::decl_iterator C) : Current(C) {
    SkipToNextDecl();
  }

  value_type operator*() const { return cast<SpecificDecl>(*Current); }
  value_type operator->() const { return cast<SpecificDecl>(*Current); }

  filtered_decl_iterator& operator++() {
    ++Current;
    SkipToNextDecl();
    return *this;
  }

  filtered_decl_iterator operator++(int) {
    filtered_decl_iterator tmp(*this);
    ++(*this);
    return tmp;
  }

  friend bool operator==(const filtered_decl_iterator& x,
                         const filtered_decl_iterator& y) {
    return x.Current == y.Current;
  }

  friend bool operator!=(const filtered_decl_iterator& x,
                         const filtered_decl_iterator& y) {
    return x.Current != y.Current;
  }
};

/// Add the declaration D into this context.
///
/// This routine should be invoked when the declaration D has first
/// been declared, to place D into the context where it was
/// (lexically) defined. Every declaration must be added to one
/// (and only one!) context, where it can be visited via
/// [decls_begin(), decls_end()). Once a declaration has been added
/// to its lexical context, the corresponding DeclContext owns the
/// declaration.
///
/// If D is also a NamedDecl, it will be made visible within its
/// semantic context via makeDeclVisibleInContext.
void addDecl(Decl *D);

/// Add the declaration D into this context, but suppress
/// searches for external declarations with the same name.
///
/// Although analogous in function to addDecl, this removes an
/// important check.  This is only useful if the Decl is being
/// added in response to an external search; in all other cases,
/// addDecl() is the right function to use.
/// See the ASTImporter for use cases.
void addDeclInternal(Decl *D);

/// Add the declaration D to this context without modifying
/// any lookup tables.
///
/// This is useful for some operations in dependent contexts where
/// the semantic context might not be dependent;  this basically
/// only happens with friends.
void addHiddenDecl(Decl *D);

/// Removes a declaration from this context.
void removeDecl(Decl *D);

/// Checks whether a declaration is in this context.
bool containsDecl(Decl *D) const;

/// Checks whether a declaration is in this context.
/// This also loads the Decls from the external source before the check.
bool containsDeclAndLoad(Decl *D) const;

using lookup_result = DeclContextLookupResult;
using lookup_iterator = lookup_result::iterator;

/// lookup - Find the declarations (if any) with the given Name in
/// this context. Returns a range of iterators that contains all of
/// the declarations with this name, with object, function, member,
/// and enumerator names preceding any tag name. Note that this
/// routine will not look into parent contexts.
lookup_result lookup(DeclarationName Name) const;

/// Find the declarations with the given name that are visible
/// within this context; don't attempt to retrieve anything from an
/// external source.
lookup_result noload_lookup(DeclarationName Name);

/// A simplistic name lookup mechanism that performs name lookup
/// into this declaration context without consulting the external source.
///
/// This function should almost never be used, because it subverts the
/// usual relationship between a DeclContext and the external source.
/// See the ASTImporter for the (few, but important) use cases.
///
/// FIXME: This is very inefficient; replace uses of it with uses of
/// noload_lookup.
void localUncachedLookup(DeclarationName Name,
                         SmallVectorImpl<NamedDecl *> &Results);

/// Makes a declaration visible within this context.
///
/// This routine makes the declaration D visible to name lookup
/// within this context and, if this is a transparent context,
/// within its parent contexts up to the first enclosing
/// non-transparent context. Making a declaration visible within a
/// context does not transfer ownership of a declaration, and a
/// declaration can be visible in many contexts that aren't its
/// lexical context.
///
/// If D is a redeclaration of an existing declaration that is
/// visible from this context, as determined by
/// NamedDecl::declarationReplaces, the previous declaration will be
/// replaced with D.
void makeDeclVisibleInContext(NamedDecl *D);

/// all_lookups_iterator - An iterator that provides a view over the results
/// of looking up every possible name.
class all_lookups_iterator;

using lookups_range = llvm::iterator_range<all_lookups_iterator>;

lookups_range lookups() const;
// Like lookups(), but avoids loading external declarations.
// If PreserveInternalState, avoids building lookup data structures too.
lookups_range noload_lookups(bool PreserveInternalState) const;

/// Iterators over all possible lookups within this context.
all_lookups_iterator lookups_begin() const;
all_lookups_iterator lookups_end() const;

/// Iterators over all possible lookups within this context that are
/// currently loaded; don't attempt to retrieve anything from an external
/// source.
all_lookups_iterator noload_lookups_begin() const;
all_lookups_iterator noload_lookups_end() const;

struct udir_iterator;

using udir_iterator_base =
    llvm::iterator_adaptor_base<udir_iterator, lookup_iterator,
                                std::random_access_iterator_tag,
                                UsingDirectiveDecl *>;

struct udir_iterator : udir_iterator_base {
  udir_iterator(lookup_iterator I) : udir_iterator_base(I) {}

  UsingDirectiveDecl *operator*() const;
};

using udir_range = llvm::iterator_range<udir_iterator>;

udir_range using_directives() const;

// These are all defined in DependentDiagnostic.h.
class ddiag_iterator;

using ddiag_range = llvm::iterator_range<DeclContext::ddiag_iterator>;

inline ddiag_range ddiags() const;

// Low-level accessors

/// Mark that there are external lexical declarations that we need
/// to include in our lookup table (and that are not available as external
/// visible lookups). These extra lookup results will be found by walking
/// the lexical declarations of this context. This should be used only if
/// setHasExternalLexicalStorage() has been called on any decl context for
/// which this is the primary context.
void setMustBuildLookupTable() {
  assert(this == getPrimaryContext() &&((this == getPrimaryContext() && "should only be called on primary context"
) ? static_cast<void> (0) : __assert_fail ("this == getPrimaryContext() && \"should only be called on primary context\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 2411, __PRETTY_FUNCTION__))
         "should only be called on primary context")((this == getPrimaryContext() && "should only be called on primary context"
) ? static_cast<void> (0) : __assert_fail ("this == getPrimaryContext() && \"should only be called on primary context\""
, "/build/llvm-toolchain-snapshot-13~++20210405022414+5f57793c4fe4/clang/include/clang/AST/DeclBase.h"
, 2411, __PRETTY_FUNCTION__));
  DeclContextBits.HasLazyExternalLexicalLookups = true;
}

/// Retrieve the internal representation of the lookup structure.
/// This may omit some names if we are lazily building the structure.
StoredDeclsMap *getLookupPtr() const { return LookupPtr; }

/// Ensure the lookup structure is fully-built and return it.
StoredDeclsMap *buildLookup();

/// Whether this DeclContext has external storage containing
/// additional declarations that are lexically in this context.
bool hasExternalLexicalStorage() const {
  return DeclContextBits.ExternalLexicalStorage;
}

/// State whether this DeclContext has external storage for
/// declarations lexically in this context.
void setHasExternalLexicalStorage(bool ES = true) const {
  DeclContextBits.ExternalLexicalStorage = ES;
}

/// Whether this DeclContext has external storage containing
/// additional declarations that are visible in this context.
bool hasExternalVisibleStorage() const {
  return DeclContextBits.ExternalVisibleStorage;
}

/// State whether this DeclContext has external storage for
/// declarations visible in this context.
void setHasExternalVisibleStorage(bool ES = true) const {
  DeclContextBits.ExternalVisibleStorage = ES;
  if (ES && LookupPtr)
    DeclContextBits.NeedToReconcileExternalVisibleStorage = true;
}

/// Determine whether the given declaration is stored in the list of
/// declarations lexically within this context.
bool isDeclInLexicalTraversal(const Decl *D) const {
  return D && (D->NextInContextAndBits.getPointer() || D == FirstDecl ||
               D == LastDecl);
}

bool setUseQualifiedLookup(bool use = true) const {
  bool old_value = DeclContextBits.UseQualifiedLookup;
  DeclContextBits.UseQualifiedLookup = use;
  return old_value;
}

bool shouldUseQualifiedLookup() const {
  return DeclContextBits.UseQualifiedLookup;
}

static bool classof(const Decl *D);
static bool classof(const DeclContext *D) { return true; }

void dumpDeclContext() const;
void dumpLookups() const;
void dumpLookups(llvm::raw_ostream &OS, bool DumpDecls = false,
                 bool Deserialize = false) const;

2473private:
/// Whether this declaration context has had externally visible
/// storage added since the last lookup. In this case, \c LookupPtr's
/// invariant may not hold and needs to be fixed before we perform
/// another lookup.
bool hasNeedToReconcileExternalVisibleStorage() const {
  return DeclContextBits.NeedToReconcileExternalVisibleStorage;
}

/// State that this declaration context has had externally visible
/// storage added since the last lookup. In this case, \c LookupPtr's
/// invariant may not hold and needs to be fixed before we perform
/// another lookup.
void setNeedToReconcileExternalVisibleStorage(bool Need = true) const {
  DeclContextBits.NeedToReconcileExternalVisibleStorage = Need;
}

/// If \c true, this context may have local lexical declarations
/// that are missing from the lookup table.
bool hasLazyLocalLexicalLookups() const {
  return DeclContextBits.HasLazyLocalLexicalLookups;
}

/// If \c true, this context may have local lexical declarations
/// that are missing from the lookup table.
void setHasLazyLocalLexicalLookups(bool HasLLLL = true) const {
  DeclContextBits.HasLazyLocalLexicalLookups = HasLLLL;
}

/// If \c true, the external source may have lexical declarations
/// that are missing from the lookup table.
bool hasLazyExternalLexicalLookups() const {
  return DeclContextBits.HasLazyExternalLexicalLookups;
}

/// If \c true, the external source may have lexical declarations
/// that are missing from the lookup table.
void setHasLazyExternalLexicalLookups(bool HasLELL = true) const {
  DeclContextBits.HasLazyExternalLexicalLookups = HasLELL;
}

void reconcileExternalVisibleStorage() const;
bool LoadLexicalDeclsFromExternalStorage() const;

/// Makes a declaration visible within this context, but
/// suppresses searches for external declarations with the same
/// name.
///
/// Analogous to makeDeclVisibleInContext, but for the exclusive
/// use of addDeclInternal().
void makeDeclVisibleInContextInternal(NamedDecl *D);

StoredDeclsMap *CreateStoredDeclsMap(ASTContext &C) const;

void loadLazyLocalLexicalLookups();
void buildLookupImpl(DeclContext *DCtx, bool Internal);
void makeDeclVisibleInContextWithFlags(NamedDecl *D, bool Internal,
                                       bool Rediscoverable);
void makeDeclVisibleInContextImpl(NamedDecl *D, bool Internal);
2532};

2534inline bool Decl::isTemplateParameter() const {
return getKind() == TemplateTypeParm || getKind() == NonTypeTemplateParm ||
       getKind() == TemplateTemplateParm;
2537}

2539// Specialization selected when ToTy is not a known subclass of DeclContext.
2540template <class ToTy,
        bool IsKnownSubtype = ::std::is_base_of<DeclContext, ToTy>::value>
2542struct cast_convert_decl_context {
static const ToTy *doit(const DeclContext *Val) {
  return static_cast<const ToTy*>(Decl::castFromDeclContext(Val));
}

static ToTy *doit(DeclContext *Val) {
  return static_cast<ToTy*>(Decl::castFromDeclContext(Val));
}
2550};

2552// Specialization selected when ToTy is a known subclass of DeclContext.
2553template <class ToTy>
2554struct cast_convert_decl_context<ToTy, true> {
static const ToTy *doit(const DeclContext *Val) {
  return static_cast<const ToTy*>(Val);
}

static ToTy *doit(DeclContext *Val) {
  return static_cast<ToTy*>(Val);
}
2562};

2564} // namespace clang

2566namespace llvm {

2568/// isa<T>(DeclContext*)
2569template <typename To>
2570struct isa_impl<To, ::clang::DeclContext> {
static bool doit(const ::clang::DeclContext &Val) {
  return To::classofKind(Val.getDeclKind());
}
2574};

2576/// cast<T>(DeclContext*)
2577template<class ToTy>
2578struct cast_convert_val<ToTy,
                      const ::clang::DeclContext,const ::clang::DeclContext> {
static const ToTy &doit(const ::clang::DeclContext &Val) {
  return *::clang::cast_convert_decl_context<ToTy>::doit(&Val);
}
2583};

2585template<class ToTy>
2586struct cast_convert_val<ToTy, ::clang::DeclContext, ::clang::DeclContext> {
static ToTy &doit(::clang::DeclContext &Val) {
  return *::clang::cast_convert_decl_context<ToTy>::doit(&Val);
}
2590};

2592template<class ToTy>
2593struct cast_convert_val<ToTy,
                   const ::clang::DeclContext*, const ::clang::DeclContext*> {
static const ToTy *doit(const ::clang::DeclContext *Val) {
  return ::clang::cast_convert_decl_context<ToTy>::doit(Val);
}
2598};

2600template<class ToTy>
2601struct cast_convert_val<ToTy, ::clang::DeclContext*, ::clang::DeclContext*> {
static ToTy *doit(::clang::DeclContext *Val) {
  return ::clang::cast_convert_decl_context<ToTy>::doit(Val);
}
2605};

2607/// Implement cast_convert_val for Decl -> DeclContext conversions.
2608template<class FromTy>
2609struct cast_convert_val< ::clang::DeclContext, FromTy, FromTy> {
static ::clang::DeclContext &doit(const FromTy &Val) {
  return *FromTy::castToDeclContext(&Val);
}
2613};

2615template<class FromTy>
2616struct cast_convert_val< ::clang::DeclContext, FromTy*, FromTy*> {
static ::clang::DeclContext *doit(const FromTy *Val) {
  return FromTy::castToDeclContext(Val);
}
2620};

2622template<class FromTy>
2623struct cast_convert_val< const ::clang::DeclContext, FromTy, FromTy> {
static const ::clang::DeclContext &doit(const FromTy &Val) {
  return *FromTy::castToDeclContext(&Val);
}
2627};

2629template<class FromTy>
2630struct cast_convert_val< const ::clang::DeclContext, FromTy*, FromTy*> {
static const ::clang::DeclContext *doit(const FromTy *Val) {
  return FromTy::castToDeclContext(Val);
}
2634};

2636} // namespace llvm

2638#endif // LLVM_CLANG_AST_DECLBASE_H