include/clang/Sema/SemaCUDA.h - llvm-project/clang - Git at Google

 //===----- SemaCUDA.h ----- Semantic Analysis for CUDA constructs ---------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 /// \file
 /// This file declares semantic analysis for CUDA constructs.
 ///
 //===----------------------------------------------------------------------===//

 #ifndef LLVM_CLANG_SEMA_SEMACUDA_H
 #define LLVM_CLANG_SEMA_SEMACUDA_H

 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/Redeclarable.h"
 #include "clang/Basic/Cuda.h"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Ownership.h"
 #include "clang/Sema/ParsedAttr.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaBase.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/SmallVector.h"
 #include <string>

 namespace clang {

 enum class CUDAFunctionTarget;

 class SemaCUDA : public SemaBase {
 public:
   SemaCUDA(Sema &S);

   /// Increments our count of the number of times we've seen a pragma forcing
   /// functions to be __host__ __device__.  So long as this count is greater
   /// than zero, all functions encountered will be __host__ __device__.
   void PushForceHostDevice();

   /// Decrements our count of the number of times we've seen a pragma forcing
   /// functions to be __host__ __device__.  Returns false if the count is 0
   /// before incrementing, so you can emit an error.
   bool PopForceHostDevice();

   ExprResult ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                  MultiExprArg ExecConfig,
                                  SourceLocation GGGLoc);

   /// A pair of a canonical FunctionDecl and a SourceLocation.  When used as the
   /// key in a hashtable, both the FD and location are hashed.
   struct FunctionDeclAndLoc {
     CanonicalDeclPtr<const FunctionDecl> FD;
     SourceLocation Loc;
   };

   /// FunctionDecls and SourceLocations for which CheckCall has emitted a
   /// (maybe deferred) "bad call" diagnostic.  We use this to avoid emitting the
   /// same deferred diag twice.
   llvm::DenseSet<FunctionDeclAndLoc> LocsWithCUDACallDiags;

   /// An inverse call graph, mapping known-emitted functions to one of their
   /// known-emitted callers (plus the location of the call).
   ///
   /// Functions that we can tell a priori must be emitted aren't added to this
   /// map.
   llvm::DenseMap</* Callee = */ CanonicalDeclPtr<const FunctionDecl>,
                  /* Caller = */ FunctionDeclAndLoc>
       DeviceKnownEmittedFns;

   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
   /// - If CurContext is a __host__ function, does not emit any diagnostics
   ///   unless \p EmitOnBothSides is true.
   /// - If CurContext is a __device__ or __global__ function, emits the
   ///   diagnostics immediately.
   /// - If CurContext is a __host__ __device__ function and we are compiling for
   ///   the device, creates a diagnostic which is emitted if and when we realize
   ///   that the function will be codegen'ed.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in CUDA device code.
   ///  if (DiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentTarget())
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
   SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);

   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as host code".
   ///
   /// Same as DiagIfDeviceCode, with "host" and "device" switched.
   SemaDiagnosticBuilder DiagIfHostCode(SourceLocation Loc, unsigned DiagID);

   /// Determines whether the given function is a CUDA device/host/kernel/etc.
   /// function.
   ///
   /// Use this rather than examining the function's attributes yourself -- you
   /// will get it wrong.  Returns CUDAFunctionTarget::Host if D is null.
   CUDAFunctionTarget IdentifyTarget(const FunctionDecl *D,
                                     bool IgnoreImplicitHDAttr = false);
   CUDAFunctionTarget IdentifyTarget(const ParsedAttributesView &Attrs);

   enum CUDAVariableTarget {
     CVT_Device,  /// Emitted on device side with a shadow variable on host side
     CVT_Host,    /// Emitted on host side only
     CVT_Both,    /// Emitted on both sides with different addresses
     CVT_Unified, /// Emitted as a unified address, e.g. managed variables
   };
   /// Determines whether the given variable is emitted on host or device side.
   CUDAVariableTarget IdentifyTarget(const VarDecl *D);

   /// Defines kinds of CUDA global host/device context where a function may be
   /// called.
   enum CUDATargetContextKind {
     CTCK_Unknown,       /// Unknown context
     CTCK_InitGlobalVar, /// Function called during global variable
                         /// initialization
   };

   /// Define the current global CUDA host/device context where a function may be
   /// called. Only used when a function is called outside of any functions.
   struct CUDATargetContext {
     CUDAFunctionTarget Target = CUDAFunctionTarget::HostDevice;
     CUDATargetContextKind Kind = CTCK_Unknown;
     Decl *D = nullptr;
   } CurCUDATargetCtx;

   struct CUDATargetContextRAII {
     SemaCUDA &S;
     SemaCUDA::CUDATargetContext SavedCtx;
     CUDATargetContextRAII(SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K,
                           Decl *D);
     ~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
   };

   /// Gets the CUDA target for the current context.
   CUDAFunctionTarget CurrentTarget() {
     return IdentifyTarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
   }

   static bool isImplicitHostDeviceFunction(const FunctionDecl *D);

   // CUDA function call preference. Must be ordered numerically from
   // worst to best.
   enum CUDAFunctionPreference {
     CFP_Never,      // Invalid caller/callee combination.
     CFP_WrongSide,  // Calls from host-device to host or device
                     // function that do not match current compilation
                     // mode.
     CFP_HostDevice, // Any calls to host/device functions.
     CFP_SameSide,   // Calls from host-device to host or device
                     // function matching current compilation mode.
     CFP_Native,     // host-to-host or device-to-device calls.
   };

   /// Identifies relative preference of a given Caller/Callee
   /// combination, based on their host/device attributes.
   /// \param Caller function which needs address of \p Callee.
   ///               nullptr in case of global context.
   /// \param Callee target function
   ///
   /// \returns preference value for particular Caller/Callee combination.
   CUDAFunctionPreference IdentifyPreference(const FunctionDecl *Caller,
                                             const FunctionDecl *Callee);

   /// Determines whether Caller may invoke Callee, based on their CUDA
   /// host/device attributes.  Returns false if the call is not allowed.
   ///
   /// Note: Will return true for CFP_WrongSide calls.  These may appear in
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
   bool IsAllowedCall(const FunctionDecl *Caller, const FunctionDecl *Callee) {
     return IdentifyPreference(Caller, Callee) != CFP_Never;
   }

   /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
   /// depending on FD and the current compilation settings.
   void maybeAddHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous);

   /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
   /// and current compilation settings.
   void MaybeAddConstantAttr(VarDecl *VD);

   /// Check whether we're allowed to call Callee from the current context.
   ///
   /// - If the call is never allowed in a semantically-correct program
   ///   (CFP_Never), emits an error and returns false.
   ///
   /// - If the call is allowed in semantically-correct programs, but only if
   ///   it's never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to
   ///   be emitted if and when the caller is codegen'ed, and returns true.
   ///
   ///   Will only create deferred diagnostics for a given SourceLocation once,
   ///   so you can safely call this multiple times without generating duplicate
   ///   deferred errors.
   ///
   /// - Otherwise, returns true without emitting any diagnostics.
   bool CheckCall(SourceLocation Loc, FunctionDecl *Callee);

   void CheckLambdaCapture(CXXMethodDecl *D, const sema::Capture &Capture);

   /// Set __device__ or __host__ __device__ attributes on the given lambda
   /// operator() method.
   ///
   /// CUDA lambdas by default is host device function unless it has explicit
   /// host or device attribute.
   void SetLambdaAttrs(CXXMethodDecl *Method);

   /// Record \p FD if it is a CUDA/HIP implicit host device function used on
   /// device side in device compilation.
   void RecordImplicitHostDeviceFuncUsedByDevice(const FunctionDecl *FD);

   /// Finds a function in \p Matches with highest calling priority
   /// from \p Caller context and erases all functions with lower
   /// calling priority.
   void EraseUnwantedMatches(
       const FunctionDecl *Caller,
       llvm::SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>>
           &Matches);

   /// Given a implicit special member, infer its CUDA target from the
   /// calls it needs to make to underlying base/field special members.
   /// \param ClassDecl the class for which the member is being created.
   /// \param CSM the kind of special member.
   /// \param MemberDecl the special member itself.
   /// \param ConstRHS true if this is a copy operation with a const object on
   ///        its RHS.
   /// \param Diagnose true if this call should emit diagnostics.
   /// \return true if there was an error inferring.
   /// The result of this call is implicit CUDA target attribute(s) attached to
   /// the member declaration.
   bool inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                            CXXSpecialMemberKind CSM,
                                            CXXMethodDecl *MemberDecl,
                                            bool ConstRHS, bool Diagnose);

   /// \return true if \p CD can be considered empty according to CUDA
   /// (E.2.3.1 in CUDA 7.5 Programming guide).
   bool isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
   bool isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *CD);

   // \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
   // case of error emits appropriate diagnostic and invalidates \p Var.
   //
   // \details CUDA allows only empty constructors as initializers for global
   // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
   // __shared__ variables whether they are local or not (they all are implicitly
   // static in CUDA). One exception is that CUDA allows constant initializers
   // for __constant__ and __device__ variables.
   void checkAllowedInitializer(VarDecl *VD);

   /// Check whether NewFD is a valid overload for CUDA. Emits
   /// diagnostics and invalidates NewFD if not.
   void checkTargetOverload(FunctionDecl *NewFD, const LookupResult &Previous);
   /// Copies target attributes from the template TD to the function FD.
   void inheritTargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD);

   /// Returns the name of the launch configuration function.  This is the name
   /// of the function that will be called to configure kernel call, with the
   /// parameters specified via <<<>>>.
   std::string getConfigureFuncName() const;

 private:
   unsigned ForceHostDeviceDepth = 0;

   friend class ASTReader;
   friend class ASTWriter;
 };

 } // namespace clang

 namespace llvm {
 // Hash a FunctionDeclAndLoc by looking at both its FunctionDecl and its
 // SourceLocation.
 template <> struct DenseMapInfo<clang::SemaCUDA::FunctionDeclAndLoc> {
   using FunctionDeclAndLoc = clang::SemaCUDA::FunctionDeclAndLoc;
   using FDBaseInfo =
       DenseMapInfo<clang::CanonicalDeclPtr<const clang::FunctionDecl>>;

   static FunctionDeclAndLoc getEmptyKey() {
     return {FDBaseInfo::getEmptyKey(), clang::SourceLocation()};
   }

   static FunctionDeclAndLoc getTombstoneKey() {
     return {FDBaseInfo::getTombstoneKey(), clang::SourceLocation()};
   }

   static unsigned getHashValue(const FunctionDeclAndLoc &FDL) {
     return hash_combine(FDBaseInfo::getHashValue(FDL.FD),
                         FDL.Loc.getHashValue());
   }

   static bool isEqual(const FunctionDeclAndLoc &LHS,
                       const FunctionDeclAndLoc &RHS) {
     return LHS.FD == RHS.FD && LHS.Loc == RHS.Loc;
   }
 };
 } // namespace llvm

 #endif // LLVM_CLANG_SEMA_SEMACUDA_H
	//===----- SemaCUDA.h ----- Semantic Analysis for CUDA constructs ---------===//
	//
	// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
	// See https://llvm.org/LICENSE.txt for license information.
	// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
	//
	//===----------------------------------------------------------------------===//
	/// \file
	/// This file declares semantic analysis for CUDA constructs.
	///
	//===----------------------------------------------------------------------===//

	#ifndef LLVM_CLANG_SEMA_SEMACUDA_H
	#define LLVM_CLANG_SEMA_SEMACUDA_H

	#include "clang/AST/Decl.h"
	#include "clang/AST/DeclCXX.h"
	#include "clang/AST/Redeclarable.h"
	#include "clang/Basic/Cuda.h"
	#include "clang/Basic/SourceLocation.h"
	#include "clang/Sema/Lookup.h"
	#include "clang/Sema/Ownership.h"
	#include "clang/Sema/ParsedAttr.h"
	#include "clang/Sema/Scope.h"
	#include "clang/Sema/ScopeInfo.h"
	#include "clang/Sema/SemaBase.h"
	#include "llvm/ADT/DenseMap.h"
	#include "llvm/ADT/SmallVector.h"
	#include <string>

	namespace clang {

	enum class CUDAFunctionTarget;

	class SemaCUDA : public SemaBase {
	public:
	SemaCUDA(Sema &S);

	/// Increments our count of the number of times we've seen a pragma forcing
	/// functions to be __host__ __device__. So long as this count is greater
	/// than zero, all functions encountered will be __host__ __device__.
	void PushForceHostDevice();

	/// Decrements our count of the number of times we've seen a pragma forcing
	/// functions to be __host__ __device__. Returns false if the count is 0
	/// before incrementing, so you can emit an error.
	bool PopForceHostDevice();

	ExprResult ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
	MultiExprArg ExecConfig,
	SourceLocation GGGLoc);

	/// A pair of a canonical FunctionDecl and a SourceLocation. When used as the
	/// key in a hashtable, both the FD and location are hashed.
	struct FunctionDeclAndLoc {
	CanonicalDeclPtr<const FunctionDecl> FD;
	SourceLocation Loc;
	};

	/// FunctionDecls and SourceLocations for which CheckCall has emitted a
	/// (maybe deferred) "bad call" diagnostic. We use this to avoid emitting the
	/// same deferred diag twice.
	llvm::DenseSet<FunctionDeclAndLoc> LocsWithCUDACallDiags;

	/// An inverse call graph, mapping known-emitted functions to one of their
	/// known-emitted callers (plus the location of the call).
	///
	/// Functions that we can tell a priori must be emitted aren't added to this
	/// map.
	llvm::DenseMap</* Callee = */ CanonicalDeclPtr<const FunctionDecl>,
	/* Caller = */ FunctionDeclAndLoc>
	DeviceKnownEmittedFns;

	/// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
	/// context is "used as device code".
	///
	/// - If CurContext is a __host__ function, does not emit any diagnostics
	/// unless \p EmitOnBothSides is true.
	/// - If CurContext is a __device__ or __global__ function, emits the
	/// diagnostics immediately.
	/// - If CurContext is a __host__ __device__ function and we are compiling for
	/// the device, creates a diagnostic which is emitted if and when we realize
	/// that the function will be codegen'ed.
	///
	/// Example usage:
	///
	/// // Variable-length arrays are not allowed in CUDA device code.
	/// if (DiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentTarget())
	/// return ExprError();
	/// // Otherwise, continue parsing as normal.
	SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);

	/// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
	/// context is "used as host code".
	///
	/// Same as DiagIfDeviceCode, with "host" and "device" switched.
	SemaDiagnosticBuilder DiagIfHostCode(SourceLocation Loc, unsigned DiagID);

	/// Determines whether the given function is a CUDA device/host/kernel/etc.
	/// function.
	///
	/// Use this rather than examining the function's attributes yourself -- you
	/// will get it wrong. Returns CUDAFunctionTarget::Host if D is null.
	CUDAFunctionTarget IdentifyTarget(const FunctionDecl *D,
	bool IgnoreImplicitHDAttr = false);
	CUDAFunctionTarget IdentifyTarget(const ParsedAttributesView &Attrs);

	enum CUDAVariableTarget {
	CVT_Device, /// Emitted on device side with a shadow variable on host side
	CVT_Host, /// Emitted on host side only
	CVT_Both, /// Emitted on both sides with different addresses
	CVT_Unified, /// Emitted as a unified address, e.g. managed variables
	};
	/// Determines whether the given variable is emitted on host or device side.
	CUDAVariableTarget IdentifyTarget(const VarDecl *D);

	/// Defines kinds of CUDA global host/device context where a function may be
	/// called.
	enum CUDATargetContextKind {
	CTCK_Unknown, /// Unknown context
	CTCK_InitGlobalVar, /// Function called during global variable
	/// initialization
	};

	/// Define the current global CUDA host/device context where a function may be
	/// called. Only used when a function is called outside of any functions.
	struct CUDATargetContext {
	CUDAFunctionTarget Target = CUDAFunctionTarget::HostDevice;
	CUDATargetContextKind Kind = CTCK_Unknown;
	Decl *D = nullptr;
	} CurCUDATargetCtx;

	struct CUDATargetContextRAII {
	SemaCUDA &S;
	SemaCUDA::CUDATargetContext SavedCtx;
	CUDATargetContextRAII(SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K,
	Decl *D);
	~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
	};

	/// Gets the CUDA target for the current context.
	CUDAFunctionTarget CurrentTarget() {
	return IdentifyTarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
	}

	static bool isImplicitHostDeviceFunction(const FunctionDecl *D);

	// CUDA function call preference. Must be ordered numerically from
	// worst to best.
	enum CUDAFunctionPreference {
	CFP_Never, // Invalid caller/callee combination.
	CFP_WrongSide, // Calls from host-device to host or device
	// function that do not match current compilation
	// mode.
	CFP_HostDevice, // Any calls to host/device functions.
	CFP_SameSide, // Calls from host-device to host or device
	// function matching current compilation mode.
	CFP_Native, // host-to-host or device-to-device calls.
	};

	/// Identifies relative preference of a given Caller/Callee
	/// combination, based on their host/device attributes.
	/// \param Caller function which needs address of \p Callee.
	/// nullptr in case of global context.
	/// \param Callee target function
	///
	/// \returns preference value for particular Caller/Callee combination.
	CUDAFunctionPreference IdentifyPreference(const FunctionDecl *Caller,
	const FunctionDecl *Callee);

	/// Determines whether Caller may invoke Callee, based on their CUDA
	/// host/device attributes. Returns false if the call is not allowed.
	///
	/// Note: Will return true for CFP_WrongSide calls. These may appear in
	/// semantically correct CUDA programs, but only if they're never codegen'ed.
	bool IsAllowedCall(const FunctionDecl Caller, const FunctionDecl Callee) {
	return IdentifyPreference(Caller, Callee) != CFP_Never;
	}

	/// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
	/// depending on FD and the current compilation settings.
	void maybeAddHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous);

	/// May add implicit CUDAConstantAttr attribute to VD, depending on VD
	/// and current compilation settings.
	void MaybeAddConstantAttr(VarDecl *VD);

	/// Check whether we're allowed to call Callee from the current context.
	///
	/// - If the call is never allowed in a semantically-correct program
	/// (CFP_Never), emits an error and returns false.
	///
	/// - If the call is allowed in semantically-correct programs, but only if
	/// it's never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to
	/// be emitted if and when the caller is codegen'ed, and returns true.
	///
	/// Will only create deferred diagnostics for a given SourceLocation once,
	/// so you can safely call this multiple times without generating duplicate
	/// deferred errors.
	///
	/// - Otherwise, returns true without emitting any diagnostics.
	bool CheckCall(SourceLocation Loc, FunctionDecl *Callee);

	void CheckLambdaCapture(CXXMethodDecl *D, const sema::Capture &Capture);

	/// Set __device__ or __host__ __device__ attributes on the given lambda
	/// operator() method.
	///
	/// CUDA lambdas by default is host device function unless it has explicit
	/// host or device attribute.
	void SetLambdaAttrs(CXXMethodDecl *Method);

	/// Record \p FD if it is a CUDA/HIP implicit host device function used on
	/// device side in device compilation.
	void RecordImplicitHostDeviceFuncUsedByDevice(const FunctionDecl *FD);

	/// Finds a function in \p Matches with highest calling priority
	/// from \p Caller context and erases all functions with lower
	/// calling priority.
	void EraseUnwantedMatches(
	const FunctionDecl *Caller,
	llvm::SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>>
	&Matches);

	/// Given a implicit special member, infer its CUDA target from the
	/// calls it needs to make to underlying base/field special members.
	/// \param ClassDecl the class for which the member is being created.
	/// \param CSM the kind of special member.
	/// \param MemberDecl the special member itself.
	/// \param ConstRHS true if this is a copy operation with a const object on
	/// its RHS.
	/// \param Diagnose true if this call should emit diagnostics.
	/// \return true if there was an error inferring.
	/// The result of this call is implicit CUDA target attribute(s) attached to
	/// the member declaration.
	bool inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
	CXXSpecialMemberKind CSM,
	CXXMethodDecl *MemberDecl,
	bool ConstRHS, bool Diagnose);

	/// \return true if \p CD can be considered empty according to CUDA
	/// (E.2.3.1 in CUDA 7.5 Programming guide).
	bool isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
	bool isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *CD);

	// \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
	// case of error emits appropriate diagnostic and invalidates \p Var.
	//
	// \details CUDA allows only empty constructors as initializers for global
	// variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
	// __shared__ variables whether they are local or not (they all are implicitly
	// static in CUDA). One exception is that CUDA allows constant initializers
	// for __constant__ and __device__ variables.
	void checkAllowedInitializer(VarDecl *VD);

	/// Check whether NewFD is a valid overload for CUDA. Emits
	/// diagnostics and invalidates NewFD if not.
	void checkTargetOverload(FunctionDecl *NewFD, const LookupResult &Previous);
	/// Copies target attributes from the template TD to the function FD.
	void inheritTargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD);

	/// Returns the name of the launch configuration function. This is the name
	/// of the function that will be called to configure kernel call, with the
	/// parameters specified via <<<>>>.
	std::string getConfigureFuncName() const;

	private:
	unsigned ForceHostDeviceDepth = 0;

	friend class ASTReader;
	friend class ASTWriter;
	};

	} // namespace clang

	namespace llvm {
	// Hash a FunctionDeclAndLoc by looking at both its FunctionDecl and its
	// SourceLocation.
	template <> struct DenseMapInfo<clang::SemaCUDA::FunctionDeclAndLoc> {
	using FunctionDeclAndLoc = clang::SemaCUDA::FunctionDeclAndLoc;
	using FDBaseInfo =
	DenseMapInfo<clang::CanonicalDeclPtr<const clang::FunctionDecl>>;

	static FunctionDeclAndLoc getEmptyKey() {
	return {FDBaseInfo::getEmptyKey(), clang::SourceLocation()};
	}

	static FunctionDeclAndLoc getTombstoneKey() {
	return {FDBaseInfo::getTombstoneKey(), clang::SourceLocation()};
	}

	static unsigned getHashValue(const FunctionDeclAndLoc &FDL) {
	return hash_combine(FDBaseInfo::getHashValue(FDL.FD),
	FDL.Loc.getHashValue());
	}

	static bool isEqual(const FunctionDeclAndLoc &LHS,
	const FunctionDeclAndLoc &RHS) {
	return LHS.FD == RHS.FD && LHS.Loc == RHS.Loc;
	}
	};
	} // namespace llvm

	#endif // LLVM_CLANG_SEMA_SEMACUDA_H