[clang] [clang] Introduce `SemaCUDA` (PR #88559)
Vlad Serebrennikov via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 12 11:55:13 PDT 2024
https://github.com/Endilll created https://github.com/llvm/llvm-project/pull/88559
This patch moves CUDA-related `Sema` function into new `SemaCUDA` class, following the recent example of SYCL, OpenACC, and HLSL. This is a part of the effort to split Sema. Additional context can be found in https://github.com/llvm/llvm-project/pull/82217, https://github.com/llvm/llvm-project/pull/84184, https://github.com/llvm/llvm-project/pull/87634.
>From 655c7c5e4b73ee832762d717bfbf566c9262a4a7 Mon Sep 17 00:00:00 2001
From: Vlad Serebrennikov <serebrennikov.vladislav at gmail.com>
Date: Fri, 12 Apr 2024 21:40:32 +0300
Subject: [PATCH 1/4] Introduce `SemaCUDA`
---
clang/include/clang/Basic/Cuda.h | 8 +
clang/include/clang/Sema/Sema.h | 301 +----------------
clang/include/clang/Sema/SemaCUDA.h | 306 ++++++++++++++++++
clang/lib/Parse/ParseDecl.cpp | 3 +-
clang/lib/Parse/ParseExpr.cpp | 3 +-
clang/lib/Parse/ParsePragma.cpp | 5 +-
clang/lib/Sema/Sema.cpp | 14 +-
clang/lib/Sema/SemaBase.cpp | 5 +-
clang/lib/Sema/SemaCUDA.cpp | 173 +++++-----
clang/lib/Sema/SemaDecl.cpp | 17 +-
clang/lib/Sema/SemaDeclAttr.cpp | 9 +-
clang/lib/Sema/SemaDeclCXX.cpp | 17 +-
clang/lib/Sema/SemaExpr.cpp | 17 +-
clang/lib/Sema/SemaExprCXX.cpp | 25 +-
clang/lib/Sema/SemaLambda.cpp | 5 +-
clang/lib/Sema/SemaOverload.cpp | 51 +--
clang/lib/Sema/SemaStmt.cpp | 5 +-
clang/lib/Sema/SemaTemplate.cpp | 11 +-
.../lib/Sema/SemaTemplateInstantiateDecl.cpp | 3 +-
clang/lib/Sema/SemaType.cpp | 9 +-
clang/lib/Serialization/ASTReader.cpp | 3 +-
clang/lib/Serialization/ASTWriter.cpp | 5 +-
22 files changed, 528 insertions(+), 467 deletions(-)
create mode 100644 clang/include/clang/Sema/SemaCUDA.h
diff --git a/clang/include/clang/Basic/Cuda.h b/clang/include/clang/Basic/Cuda.h
index 3e77a74c7c0092..acc6bb6581d857 100644
--- a/clang/include/clang/Basic/Cuda.h
+++ b/clang/include/clang/Basic/Cuda.h
@@ -126,6 +126,14 @@ enum class CudaArch {
HIPDefault = CudaArch::GFX906,
};
+enum class CUDAFunctionTarget {
+ Device,
+ Global,
+ Host,
+ HostDevice,
+ InvalidTarget
+};
+
static inline bool IsNVIDIAGpuArch(CudaArch A) {
return A >= CudaArch::SM_20 && A < CudaArch::GFX600;
}
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 00888b7f7a738e..5dac8083714b61 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -38,6 +38,7 @@
#include "clang/AST/TypeOrdering.h"
#include "clang/Basic/BitmaskEnum.h"
#include "clang/Basic/Builtins.h"
+#include "clang/Basic/Cuda.h"
#include "clang/Basic/DarwinSDKInfo.h"
#include "clang/Basic/ExpressionTraits.h"
#include "clang/Basic/Module.h"
@@ -183,6 +184,7 @@ class Preprocessor;
class PseudoDestructorTypeStorage;
class PseudoObjectExpr;
class QualType;
+class SemaCUDA;
class SemaHLSL;
class SemaOpenACC;
class SemaSYCL;
@@ -435,14 +437,6 @@ enum class CXXSpecialMemberKind {
Invalid
};
-enum class CUDAFunctionTarget {
- Device,
- Global,
- Host,
- HostDevice,
- InvalidTarget
-};
-
/// Sema - This implements semantic analysis and AST building for C.
/// \nosubgrouping
class Sema final : public SemaBase {
@@ -981,9 +975,19 @@ class Sema final : public SemaBase {
return DelayedDiagnostics.push(pool);
}
+ /// Diagnostics that are emitted only if we discover that the given function
+ /// must be codegen'ed. Because handling these correctly adds overhead to
+ /// compilation, this is currently only enabled for CUDA compilations.
+ SemaDiagnosticBuilder::DeferredDiagnosticsType DeviceDeferredDiags;
+
/// CurContext - This is the current declaration context of parsing.
DeclContext *CurContext;
+ SemaCUDA &CUDA() {
+ assert(CUDAPtr);
+ return *CUDAPtr;
+ }
+
SemaHLSL &HLSL() {
assert(HLSLPtr);
return *HLSLPtr;
@@ -1029,6 +1033,7 @@ class Sema final : public SemaBase {
mutable IdentifierInfo *Ident_super;
+ std::unique_ptr<SemaCUDA> CUDAPtr;
std::unique_ptr<SemaHLSL> HLSLPtr;
std::unique_ptr<SemaOpenACC> OpenACCPtr;
std::unique_ptr<SemaSYCL> SYCLPtr;
@@ -12908,258 +12913,6 @@ class Sema final : public SemaBase {
//
//
- /// \name CUDA
- /// Implementations are in SemaCUDA.cpp
- ///@{
-
-public:
- /// 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 PushForceCUDAHostDevice();
-
- /// 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 PopForceCUDAHostDevice();
-
- ExprResult ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
- MultiExprArg ExecConfig,
- SourceLocation GGGLoc);
-
- /// Diagnostics that are emitted only if we discover that the given function
- /// must be codegen'ed. Because handling these correctly adds overhead to
- /// compilation, this is currently only enabled for CUDA compilations.
- SemaDiagnosticBuilder::DeferredDiagnosticsType DeviceDeferredDiags;
-
- /// A pair of a canonical FunctionDecl and a SourceLocation. When used as the
- /// key in a hashtable, both the FD and location are hashed.
- struct FunctionDeclAndLoc {
- CanonicalDeclPtr<const FunctionDecl> FD;
- SourceLocation Loc;
- };
-
- /// FunctionDecls and SourceLocations for which CheckCUDACall has emitted a
- /// (maybe deferred) "bad call" diagnostic. We use this to avoid emitting the
- /// same deferred diag twice.
- llvm::DenseSet<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 (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla)
- /// << llvm::to_underlying(CurrentCUDATarget()))
- /// return ExprError();
- /// // Otherwise, continue parsing as normal.
- SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
- unsigned DiagID);
-
- /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
- /// context is "used as host code".
- ///
- /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
- SemaDiagnosticBuilder CUDADiagIfHostCode(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 IdentifyCUDATarget(const FunctionDecl *D,
- bool IgnoreImplicitHDAttr = false);
- CUDAFunctionTarget IdentifyCUDATarget(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 IdentifyCUDATarget(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 {
- Sema &S;
- CUDATargetContext SavedCtx;
- CUDATargetContextRAII(Sema &S_, CUDATargetContextKind K, Decl *D);
- ~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
- };
-
- /// Gets the CUDA target for the current context.
- CUDAFunctionTarget CurrentCUDATarget() {
- return IdentifyCUDATarget(dyn_cast<FunctionDecl>(CurContext));
- }
-
- static bool isCUDAImplicitHostDeviceFunction(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 IdentifyCUDAPreference(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 IsAllowedCUDACall(const FunctionDecl *Caller,
- const FunctionDecl *Callee) {
- return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
- }
-
- /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
- /// depending on FD and the current compilation settings.
- void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
- const LookupResult &Previous);
-
- /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
- /// and current compilation settings.
- void MaybeAddCUDAConstantAttr(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 CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
-
- void CUDACheckLambdaCapture(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 CUDASetLambdaAttrs(CXXMethodDecl *Method);
-
- /// Record \p FD if it is a CUDA/HIP implicit host device function used on
- /// device side in device compilation.
- void CUDARecordImplicitHostDeviceFuncUsedByDevice(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 EraseUnwantedCUDAMatches(
- const FunctionDecl *Caller,
- 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 inferCUDATargetForImplicitSpecialMember(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 isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
- bool isEmptyCudaDestructor(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 checkAllowedCUDAInitializer(VarDecl *VD);
-
- /// Check whether NewFD is a valid overload for CUDA. Emits
- /// diagnostics and invalidates NewFD if not.
- void checkCUDATargetOverload(FunctionDecl *NewFD,
- const LookupResult &Previous);
- /// Copies target attributes from the template TD to the function FD.
- void inheritCUDATargetAttrs(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 getCudaConfigureFuncName() const;
-
-private:
- unsigned ForceCUDAHostDeviceDepth = 0;
-
- ///@}
-
- //
- //
- // -------------------------------------------------------------------------
- //
- //
-
/// \name OpenMP Directives and Clauses
/// Implementations are in SemaOpenMP.cpp
///@{
@@ -14546,32 +14299,4 @@ std::unique_ptr<sema::RISCVIntrinsicManager>
CreateRISCVIntrinsicManager(Sema &S);
} // end namespace clang
-namespace llvm {
-// Hash a FunctionDeclAndLoc by looking at both its FunctionDecl and its
-// SourceLocation.
-template <> struct DenseMapInfo<clang::Sema::FunctionDeclAndLoc> {
- using FunctionDeclAndLoc = clang::Sema::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
diff --git a/clang/include/clang/Sema/SemaCUDA.h b/clang/include/clang/Sema/SemaCUDA.h
new file mode 100644
index 00000000000000..1ceaab06a5cc62
--- /dev/null
+++ b/clang/include/clang/Sema/SemaCUDA.h
@@ -0,0 +1,306 @@
+//===----- 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 PushForceCUDAHostDevice();
+
+ /// 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 PopForceCUDAHostDevice();
+
+ ExprResult ActOnCUDAExecConfigExpr(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 CheckCUDACall 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 (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
+ /// return ExprError();
+ /// // Otherwise, continue parsing as normal.
+ SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
+ unsigned DiagID);
+
+ /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+ /// context is "used as host code".
+ ///
+ /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
+ SemaDiagnosticBuilder CUDADiagIfHostCode(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 IdentifyCUDATarget(const FunctionDecl *D,
+ bool IgnoreImplicitHDAttr = false);
+ CUDAFunctionTarget IdentifyCUDATarget(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 IdentifyCUDATarget(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 CurrentCUDATarget() {
+ return IdentifyCUDATarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
+ }
+
+ static bool isCUDAImplicitHostDeviceFunction(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 IdentifyCUDAPreference(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 IsAllowedCUDACall(const FunctionDecl *Caller,
+ const FunctionDecl *Callee) {
+ return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
+ }
+
+ /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
+ /// depending on FD and the current compilation settings.
+ void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
+ const LookupResult &Previous);
+
+ /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
+ /// and current compilation settings.
+ void MaybeAddCUDAConstantAttr(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 CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
+
+ void CUDACheckLambdaCapture(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 CUDASetLambdaAttrs(CXXMethodDecl *Method);
+
+ /// Record \p FD if it is a CUDA/HIP implicit host device function used on
+ /// device side in device compilation.
+ void CUDARecordImplicitHostDeviceFuncUsedByDevice(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 EraseUnwantedCUDAMatches(
+ 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 inferCUDATargetForImplicitSpecialMember(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 isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
+ bool isEmptyCudaDestructor(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 checkAllowedCUDAInitializer(VarDecl *VD);
+
+ /// Check whether NewFD is a valid overload for CUDA. Emits
+ /// diagnostics and invalidates NewFD if not.
+ void checkCUDATargetOverload(FunctionDecl *NewFD,
+ const LookupResult &Previous);
+ /// Copies target attributes from the template TD to the function FD.
+ void inheritCUDATargetAttrs(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 getCudaConfigureFuncName() const;
+
+private:
+ unsigned ForceCUDAHostDeviceDepth = 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
\ No newline at end of file
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 583232f2d610d0..37be68d6ec5be2 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -26,6 +26,7 @@
#include "clang/Sema/Lookup.h"
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaDiagnostic.h"
#include "llvm/ADT/SmallSet.h"
#include "llvm/ADT/SmallString.h"
@@ -2664,7 +2665,7 @@ Decl *Parser::ParseDeclarationAfterDeclaratorAndAttributes(
}
}
- Sema::CUDATargetContextRAII X(Actions, Sema::CTCK_InitGlobalVar, ThisDecl);
+ SemaCUDA::CUDATargetContextRAII X(Actions.CUDA(), SemaCUDA::CTCK_InitGlobalVar, ThisDecl);
switch (TheInitKind) {
// Parse declarator '=' initializer.
case InitKind::Equal: {
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index d08e675604d19c..e2410555f5e0de 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -30,6 +30,7 @@
#include "clang/Sema/EnterExpressionEvaluationContext.h"
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaSYCL.h"
#include "clang/Sema/TypoCorrection.h"
#include "llvm/ADT/SmallVector.h"
@@ -2129,7 +2130,7 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
}
if (!LHS.isInvalid()) {
- ExprResult ECResult = Actions.ActOnCUDAExecConfigExpr(getCurScope(),
+ ExprResult ECResult = Actions.CUDA().ActOnCUDAExecConfigExpr(getCurScope(),
OpenLoc,
ExecConfigExprs,
CloseLoc);
diff --git a/clang/lib/Parse/ParsePragma.cpp b/clang/lib/Parse/ParsePragma.cpp
index 0f692e2146a490..cd784eb0fb0cdc 100644
--- a/clang/lib/Parse/ParsePragma.cpp
+++ b/clang/lib/Parse/ParsePragma.cpp
@@ -21,6 +21,7 @@
#include "clang/Parse/RAIIObjectsForParser.h"
#include "clang/Sema/EnterExpressionEvaluationContext.h"
#include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/StringSwitch.h"
#include <optional>
@@ -3900,8 +3901,8 @@ void PragmaForceCUDAHostDeviceHandler::HandlePragma(
}
if (Info->isStr("begin"))
- Actions.PushForceCUDAHostDevice();
- else if (!Actions.PopForceCUDAHostDevice())
+ Actions.CUDA().PushForceCUDAHostDevice();
+ else if (!Actions.CUDA().PopForceCUDAHostDevice())
PP.Diag(FirstTok.getLocation(),
diag::err_pragma_cannot_end_force_cuda_host_device);
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index a2ea66f339c8e3..50de4f33ba8b69 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -41,6 +41,7 @@
#include "clang/Sema/RISCVIntrinsicManager.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaConsumer.h"
#include "clang/Sema/SemaHLSL.h"
#include "clang/Sema/SemaInternal.h"
@@ -199,6 +200,7 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer,
LateTemplateParser(nullptr), LateTemplateParserCleanup(nullptr),
OpaqueParser(nullptr), CurContext(nullptr), ExternalSource(nullptr),
CurScope(nullptr), Ident_super(nullptr),
+ CUDAPtr(std::make_unique<SemaCUDA>(*this)),
HLSLPtr(std::make_unique<SemaHLSL>(*this)),
OpenACCPtr(std::make_unique<SemaOpenACC>(*this)),
SYCLPtr(std::make_unique<SemaSYCL>(*this)),
@@ -1635,15 +1637,15 @@ bool Sema::hasUncompilableErrorOccurred() const {
// Print notes showing how we can reach FD starting from an a priori
// known-callable function.
static void emitCallStackNotes(Sema &S, const FunctionDecl *FD) {
- auto FnIt = S.DeviceKnownEmittedFns.find(FD);
- while (FnIt != S.DeviceKnownEmittedFns.end()) {
+ auto FnIt = S.CUDA().DeviceKnownEmittedFns.find(FD);
+ while (FnIt != S.CUDA().DeviceKnownEmittedFns.end()) {
// Respect error limit.
if (S.Diags.hasFatalErrorOccurred())
return;
DiagnosticBuilder Builder(
S.Diags.Report(FnIt->second.Loc, diag::note_called_by));
Builder << FnIt->second.FD;
- FnIt = S.DeviceKnownEmittedFns.find(FnIt->second.FD);
+ FnIt = S.CUDA().DeviceKnownEmittedFns.find(FnIt->second.FD);
}
}
@@ -1747,7 +1749,7 @@ class DeferredDiagnosticsEmitter
(ShouldEmitRootNode || InOMPDeviceContext))
S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc);
if (Caller)
- S.DeviceKnownEmittedFns[FD] = {Caller, Loc};
+ S.CUDA().DeviceKnownEmittedFns[FD] = {Caller, Loc};
// Always emit deferred diagnostics for the direct users. This does not
// lead to explosion of diagnostics since each user is visited at most
// twice.
@@ -1900,8 +1902,8 @@ Sema::targetDiag(SourceLocation Loc, unsigned DiagID, const FunctionDecl *FD) {
? diagIfOpenMPDeviceCode(Loc, DiagID, FD)
: diagIfOpenMPHostCode(Loc, DiagID, FD);
if (getLangOpts().CUDA)
- return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
- : CUDADiagIfHostCode(Loc, DiagID);
+ return getLangOpts().CUDAIsDevice ? CUDA().CUDADiagIfDeviceCode(Loc, DiagID)
+ : CUDA().CUDADiagIfHostCode(Loc, DiagID);
if (getLangOpts().SYCLIsDevice)
return SYCL().DiagIfDeviceCode(Loc, DiagID);
diff --git a/clang/lib/Sema/SemaBase.cpp b/clang/lib/Sema/SemaBase.cpp
index 95c0cfbe283b0e..17cff0e84bb0d3 100644
--- a/clang/lib/Sema/SemaBase.cpp
+++ b/clang/lib/Sema/SemaBase.cpp
@@ -1,5 +1,6 @@
#include "clang/Sema/SemaBase.h"
#include "clang/Sema/Sema.h"
+#include "clang/Sema/SemaCUDA.h"
namespace clang {
@@ -70,8 +71,8 @@ Sema::SemaDiagnosticBuilder SemaBase::Diag(SourceLocation Loc, unsigned DiagID,
}
SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice
- ? SemaRef.CUDADiagIfDeviceCode(Loc, DiagID)
- : SemaRef.CUDADiagIfHostCode(Loc, DiagID);
+ ? SemaRef.CUDA().CUDADiagIfDeviceCode(Loc, DiagID)
+ : SemaRef.CUDA().CUDADiagIfHostCode(Loc, DiagID);
SetIsLastErrorImmediate(DB.isImmediate());
return DB;
}
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 9d6d709e262ad1..6e94e4e8091f60 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -10,6 +10,7 @@
///
//===----------------------------------------------------------------------===//
+#include "clang/Sema/SemaCUDA.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Decl.h"
#include "clang/AST/ExprCXX.h"
@@ -27,6 +28,8 @@
#include <optional>
using namespace clang;
+SemaCUDA::SemaCUDA(Sema &S) : SemaBase(S) {}
+
template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) {
if (!D)
return false;
@@ -35,12 +38,12 @@ template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) {
return false;
}
-void Sema::PushForceCUDAHostDevice() {
+void SemaCUDA::PushForceCUDAHostDevice() {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
ForceCUDAHostDeviceDepth++;
}
-bool Sema::PopForceCUDAHostDevice() {
+bool SemaCUDA::PopForceCUDAHostDevice() {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
if (ForceCUDAHostDeviceDepth == 0)
return false;
@@ -48,24 +51,24 @@ bool Sema::PopForceCUDAHostDevice() {
return true;
}
-ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
+ExprResult SemaCUDA::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc) {
- FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
+ FunctionDecl *ConfigDecl = getASTContext().getcudaConfigureCallDecl();
if (!ConfigDecl)
return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
<< getCudaConfigureFuncName());
QualType ConfigQTy = ConfigDecl->getType();
- DeclRefExpr *ConfigDR = new (Context)
- DeclRefExpr(Context, ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
- MarkFunctionReferenced(LLLLoc, ConfigDecl);
+ DeclRefExpr *ConfigDR = new (getASTContext())
+ DeclRefExpr(getASTContext(), ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
+ SemaRef.MarkFunctionReferenced(LLLLoc, ConfigDecl);
- return BuildCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
+ return SemaRef.BuildCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
/*IsExecConfig=*/true);
}
-CUDAFunctionTarget Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
+CUDAFunctionTarget SemaCUDA::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
bool HasHostAttr = false;
bool HasDeviceAttr = false;
bool HasGlobalAttr = false;
@@ -112,12 +115,12 @@ static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
});
}
-Sema::CUDATargetContextRAII::CUDATargetContextRAII(Sema &S_,
- CUDATargetContextKind K,
+SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(SemaCUDA &S_,
+ SemaCUDA::CUDATargetContextKind K,
Decl *D)
: S(S_) {
SavedCtx = S.CurCUDATargetCtx;
- assert(K == CTCK_InitGlobalVar);
+ assert(K == SemaCUDA::CTCK_InitGlobalVar);
auto *VD = dyn_cast_or_null<VarDecl>(D);
if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
auto Target = CUDAFunctionTarget::Host;
@@ -131,7 +134,7 @@ Sema::CUDATargetContextRAII::CUDATargetContextRAII(Sema &S_,
}
/// IdentifyCUDATarget - Determine the CUDA compilation target for this function
-CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
+CUDAFunctionTarget SemaCUDA::IdentifyCUDATarget(const FunctionDecl *D,
bool IgnoreImplicitHDAttr) {
// Code that lives outside a function gets the target from CurCUDATargetCtx.
if (D == nullptr)
@@ -160,7 +163,7 @@ CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
}
/// IdentifyTarget - Determine the CUDA compilation target for this variable.
-Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
+SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyCUDATarget(const VarDecl *Var) {
if (Var->hasAttr<HIPManagedAttr>())
return CVT_Unified;
// Only constexpr and const variabless with implicit constant attribute
@@ -221,8 +224,8 @@ Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
// | hd | h | SS | WS | (d) |
// | hd | hd | HD | HD | (b) |
-Sema::CUDAFunctionPreference
-Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
+SemaCUDA::CUDAFunctionPreference
+SemaCUDA::IdentifyCUDAPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
assert(Callee && "Callee must be valid.");
@@ -309,13 +312,13 @@ template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) {
return D->isImplicit();
}
-bool Sema::isCUDAImplicitHostDeviceFunction(const FunctionDecl *D) {
+bool SemaCUDA::isCUDAImplicitHostDeviceFunction(const FunctionDecl *D) {
bool IsImplicitDevAttr = hasImplicitAttr<CUDADeviceAttr>(D);
bool IsImplicitHostAttr = hasImplicitAttr<CUDAHostAttr>(D);
return IsImplicitDevAttr && IsImplicitHostAttr;
}
-void Sema::EraseUnwantedCUDAMatches(
+void SemaCUDA::EraseUnwantedCUDAMatches(
const FunctionDecl *Caller,
SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
if (Matches.size() <= 1)
@@ -367,7 +370,7 @@ resolveCalleeCUDATargetConflict(CUDAFunctionTarget Target1,
return false;
}
-bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
+bool SemaCUDA::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXSpecialMemberKind CSM,
CXXMethodDecl *MemberDecl,
bool ConstRHS,
@@ -388,7 +391,7 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
// We're going to invoke special member lookup; mark that these special
// members are called from this one, and not from its caller.
- ContextRAII MethodContext(*this, MemberDecl);
+ Sema::ContextRAII MethodContext(SemaRef, MemberDecl);
// Look for special members in base classes that should be invoked from here.
// Infer the target of this member base on the ones it should call.
@@ -412,7 +415,7 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
Sema::SpecialMemberOverloadResult SMOR =
- LookupSpecialMember(BaseClassDecl, CSM,
+ SemaRef.LookupSpecialMember(BaseClassDecl, CSM,
/* ConstArg */ ConstRHS,
/* VolatileArg */ false,
/* RValueThis */ false,
@@ -435,7 +438,7 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
<< (unsigned)CSM << llvm::to_underlying(*InferredTarget)
<< llvm::to_underlying(BaseMethodTarget);
}
- MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
+ MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
return true;
}
}
@@ -448,14 +451,14 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
}
const RecordType *FieldType =
- Context.getBaseElementType(F->getType())->getAs<RecordType>();
+ getASTContext().getBaseElementType(F->getType())->getAs<RecordType>();
if (!FieldType) {
continue;
}
CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
Sema::SpecialMemberOverloadResult SMOR =
- LookupSpecialMember(FieldRecDecl, CSM,
+ SemaRef.LookupSpecialMember(FieldRecDecl, CSM,
/* ConstArg */ ConstRHS && !F->isMutable(),
/* VolatileArg */ false,
/* RValueThis */ false,
@@ -479,7 +482,7 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
<< (unsigned)CSM << llvm::to_underlying(*InferredTarget)
<< llvm::to_underlying(FieldMethodTarget);
}
- MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
+ MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
return true;
}
}
@@ -499,16 +502,16 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
// We either setting attributes first time, or the inferred ones must match
// previously set ones.
if (NeedsD && !HasD)
- MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext()));
if (NeedsH && !HasH)
- MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
return false;
}
-bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
+bool SemaCUDA::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
if (!CD->isDefined() && CD->isTemplateInstantiation())
- InstantiateFunctionDefinition(Loc, CD->getFirstDecl());
+ SemaRef.InstantiateFunctionDefinition(Loc, CD->getFirstDecl());
// (E.2.3.1, CUDA 7.5) A constructor for a class type is considered
// empty at a point in the translation unit, if it is either a
@@ -544,13 +547,13 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
return true;
}
-bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
+bool SemaCUDA::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
// No destructor -> no problem.
if (!DD)
return true;
if (!DD->isDefined() && DD->isTemplateInstantiation())
- InstantiateFunctionDefinition(Loc, DD->getFirstDecl());
+ SemaRef.InstantiateFunctionDefinition(Loc, DD->getFirstDecl());
// (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
// empty at a point in the translation unit, if it is either a
@@ -620,7 +623,7 @@ bool IsDependentVar(VarDecl *VD) {
// __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.
-bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
+bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
CUDAInitializerCheckKind CheckKind) {
assert(!VD->isInvalidDecl() && VD->hasGlobalStorage());
assert(!IsDependentVar(VD) && "do not check dependent var");
@@ -635,9 +638,9 @@ bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
};
auto IsConstantInit = [&](const Expr *Init) {
assert(Init);
- ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.Context,
+ ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.getASTContext(),
/*NoWronSidedVars=*/true);
- return Init->isConstantInitializer(S.Context,
+ return Init->isConstantInitializer(S.getASTContext(),
VD->getType()->isReferenceType());
};
auto HasEmptyDtor = [&](VarDecl *VD) {
@@ -647,12 +650,12 @@ bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
};
if (CheckKind == CICK_Shared)
return IsEmptyInit(Init) && HasEmptyDtor(VD);
- return S.LangOpts.GPUAllowDeviceInit ||
+ return S.getLangOpts().GPUAllowDeviceInit ||
((IsEmptyInit(Init) || IsConstantInit(Init)) && HasEmptyDtor(VD));
}
} // namespace
-void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
+void SemaCUDA::checkAllowedCUDAInitializer(VarDecl *VD) {
// Return early if VD is inside a non-instantiated template function since
// the implicit constructor is not defined yet.
if (const FunctionDecl *FD =
@@ -700,9 +703,9 @@ void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
}
}
-void Sema::CUDARecordImplicitHostDeviceFuncUsedByDevice(
+void SemaCUDA::CUDARecordImplicitHostDeviceFuncUsedByDevice(
const FunctionDecl *Callee) {
- FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
if (!Caller)
return;
@@ -734,15 +737,15 @@ void Sema::CUDARecordImplicitHostDeviceFuncUsedByDevice(
// ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
// #pragma clang force_cuda_host_device_begin/end
// pair).
-void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
+void SemaCUDA::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
if (ForceCUDAHostDeviceDepth > 0) {
if (!NewD->hasAttr<CUDAHostAttr>())
- NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ NewD->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
if (!NewD->hasAttr<CUDADeviceAttr>())
- NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ NewD->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext()));
return;
}
@@ -753,8 +756,8 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
!NewD->hasAttr<CUDAGlobalAttr>() &&
(NewD->getDescribedFunctionTemplate() ||
NewD->isFunctionTemplateSpecialization())) {
- NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
- NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ NewD->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
+ NewD->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext()));
return;
}
@@ -771,7 +774,7 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
FunctionDecl *OldD = D->getAsFunction();
return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
!OldD->hasAttr<CUDAHostAttr>() &&
- !IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
+ !SemaRef.IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
/* ConsiderCudaAttrs = */ false);
};
auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
@@ -781,7 +784,7 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
// in a system header, in which case we simply return without making NewD
// host+device.
NamedDecl *Match = *It;
- if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
+ if (!SemaRef.getSourceManager().isInSystemHeader(Match->getLocation())) {
Diag(NewD->getLocation(),
diag::err_cuda_unattributed_constexpr_cannot_overload_device)
<< NewD;
@@ -791,14 +794,14 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
return;
}
- NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
- NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ NewD->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
+ NewD->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext()));
}
// TODO: `__constant__` memory may be a limited resource for certain targets.
// A safeguard may be needed at the end of compilation pipeline if
// `__constant__` memory usage goes beyond limit.
-void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
+void SemaCUDA::MaybeAddCUDAConstantAttr(VarDecl *VD) {
// Do not promote dependent variables since the cotr/dtor/initializer are
// not determined. Do it after instantiation.
if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
@@ -812,10 +815,10 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
}
}
-Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
+SemaBase::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfDeviceCode(SourceLocation Loc,
unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
@@ -829,23 +832,23 @@ Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
// mode until the function is known-emitted.
if (!getLangOpts().CUDAIsDevice)
return SemaDiagnosticBuilder::K_Nop;
- if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
+ if (SemaRef.IsLastErrorImmediate && getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
return SemaDiagnosticBuilder::K_Immediate;
- return (getEmissionStatus(CurFunContext) ==
- FunctionEmissionStatus::Emitted)
+ return (SemaRef.getEmissionStatus(CurFunContext) ==
+ Sema::FunctionEmissionStatus::Emitted)
? SemaDiagnosticBuilder::K_ImmediateWithCallStack
: SemaDiagnosticBuilder::K_Deferred;
default:
return SemaDiagnosticBuilder::K_Nop;
}
}();
- return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
+ return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
}
-Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
+Sema::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfHostCode(SourceLocation Loc,
unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
@@ -858,37 +861,37 @@ Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
// mode until the function is known-emitted.
if (getLangOpts().CUDAIsDevice)
return SemaDiagnosticBuilder::K_Nop;
- if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
+ if (SemaRef.IsLastErrorImmediate && getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
return SemaDiagnosticBuilder::K_Immediate;
- return (getEmissionStatus(CurFunContext) ==
- FunctionEmissionStatus::Emitted)
+ return (SemaRef.getEmissionStatus(CurFunContext) ==
+ Sema::FunctionEmissionStatus::Emitted)
? SemaDiagnosticBuilder::K_ImmediateWithCallStack
: SemaDiagnosticBuilder::K_Deferred;
default:
return SemaDiagnosticBuilder::K_Nop;
}
}();
- return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
+ return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
}
-bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
+bool SemaCUDA::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
assert(Callee && "Callee may not be null.");
- const auto &ExprEvalCtx = currentEvaluationContext();
+ const auto &ExprEvalCtx = SemaRef.currentEvaluationContext();
if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
return true;
// FIXME: Is bailing out early correct here? Should we instead assume that
// the caller is a global initializer?
- FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
if (!Caller)
return true;
// If the caller is known-emitted, mark the callee as known-emitted.
// Otherwise, mark the call in our call graph so we can traverse it later.
bool CallerKnownEmitted =
- getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
+ SemaRef.getEmissionStatus(Caller) == Sema::FunctionEmissionStatus::Emitted;
SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
CallerKnownEmitted] {
switch (IdentifyCUDAPreference(Caller, Callee)) {
@@ -908,7 +911,7 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
if (DiagKind == SemaDiagnosticBuilder::K_Nop) {
// For -fgpu-rdc, keep track of external kernels used by host functions.
- if (LangOpts.CUDAIsDevice && LangOpts.GPURelocatableDeviceCode &&
+ if (getLangOpts().CUDAIsDevice && getLangOpts().GPURelocatableDeviceCode &&
Callee->hasAttr<CUDAGlobalAttr>() && !Callee->isDefined() &&
(!Caller || (!Caller->getDescribedFunctionTemplate() &&
getASTContext().GetGVALinkageForFunction(Caller) ==
@@ -924,12 +927,12 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
return true;
- SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+ SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, SemaRef)
<< llvm::to_underlying(IdentifyCUDATarget(Callee)) << /*function*/ 0
<< Callee << llvm::to_underlying(IdentifyCUDATarget(Caller));
if (!Callee->getBuiltinID())
SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
- diag::note_previous_decl, Caller, *this)
+ diag::note_previous_decl, Caller, SemaRef)
<< Callee;
return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
@@ -940,7 +943,7 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
// defined and uses the capture by reference when the lambda is called. When
// the capture and use happen on different sides, the capture is invalid and
// should be diagnosed.
-void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
+void SemaCUDA::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
const sema::Capture &Capture) {
// In host compilation we only need to check lambda functions emitted on host
// side. In such lambda functions, a reference capture is invalid only
@@ -950,12 +953,12 @@ void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
// kernel cannot pass a lambda back to a host function since we cannot
// define a kernel argument type which can hold the lambda before the lambda
// itself is defined.
- if (!LangOpts.CUDAIsDevice)
+ if (!getLangOpts().CUDAIsDevice)
return;
// File-scope lambda can only do init captures for global variables, which
// results in passing by value for these global variables.
- FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
if (!Caller)
return;
@@ -972,7 +975,7 @@ void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
if (Capture.isVariableCapture() && !getLangOpts().HIPStdPar) {
SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
- diag::err_capture_bad_target, Callee, *this)
+ diag::err_capture_bad_target, Callee, SemaRef)
<< Capture.getVariable();
} else if (Capture.isThisCapture()) {
// Capture of this pointer is allowed since this pointer may be pointing to
@@ -981,19 +984,19 @@ void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
// accessible on device side.
SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
diag::warn_maybe_capture_bad_target_this_ptr, Callee,
- *this);
+ SemaRef);
}
}
-void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
+void SemaCUDA::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
return;
- Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
- Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ Method->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext()));
+ Method->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
}
-void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
+void SemaCUDA::checkCUDATargetOverload(FunctionDecl *NewFD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
CUDAFunctionTarget NewTarget = IdentifyCUDATarget(NewFD);
@@ -1010,16 +1013,16 @@ void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
// should have the same implementation on both sides.
if (NewTarget != OldTarget &&
((NewTarget == CUDAFunctionTarget::HostDevice &&
- !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+ !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isCUDAImplicitHostDeviceFunction(NewFD) &&
OldTarget == CUDAFunctionTarget::Device)) ||
(OldTarget == CUDAFunctionTarget::HostDevice &&
- !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+ !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isCUDAImplicitHostDeviceFunction(OldFD) &&
NewTarget == CUDAFunctionTarget::Device)) ||
(NewTarget == CUDAFunctionTarget::Global) ||
(OldTarget == CUDAFunctionTarget::Global)) &&
- !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
+ !SemaRef.IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
/* ConsiderCudaAttrs = */ false)) {
Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
<< llvm::to_underlying(NewTarget) << NewFD->getDeclName()
@@ -1041,21 +1044,21 @@ static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
}
}
-void Sema::inheritCUDATargetAttrs(FunctionDecl *FD,
+void SemaCUDA::inheritCUDATargetAttrs(FunctionDecl *FD,
const FunctionTemplateDecl &TD) {
const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
- copyAttrIfPresent<CUDAGlobalAttr>(*this, FD, TemplateFD);
- copyAttrIfPresent<CUDAHostAttr>(*this, FD, TemplateFD);
- copyAttrIfPresent<CUDADeviceAttr>(*this, FD, TemplateFD);
+ copyAttrIfPresent<CUDAGlobalAttr>(SemaRef, FD, TemplateFD);
+ copyAttrIfPresent<CUDAHostAttr>(SemaRef, FD, TemplateFD);
+ copyAttrIfPresent<CUDADeviceAttr>(SemaRef, FD, TemplateFD);
}
-std::string Sema::getCudaConfigureFuncName() const {
+std::string SemaCUDA::getCudaConfigureFuncName() const {
if (getLangOpts().HIP)
return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration"
: "hipConfigureCall";
// New CUDA kernel launch sequence.
- if (CudaFeatureEnabled(Context.getTargetInfo().getSDKVersion(),
+ if (CudaFeatureEnabled(getASTContext().getTargetInfo().getSDKVersion(),
CudaFeature::CUDA_USES_NEW_LAUNCH))
return "__cudaPushCallConfiguration";
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 720e56692359b3..7d9efb824a96d4 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -45,6 +45,7 @@
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaHLSL.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
@@ -10600,7 +10601,7 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
// specialization inherits its target attributes from its template
// in the CheckFunctionTemplateSpecialization() call below.
if (getLangOpts().CUDA && !isFunctionTemplateSpecialization)
- maybeAddCUDAHostDeviceAttrs(NewFD, Previous);
+ CUDA().maybeAddCUDAHostDeviceAttrs(NewFD, Previous);
// Handle explict specializations of function templates
// and friend function declarations with an explicit
@@ -10898,12 +10899,12 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
if (getLangOpts().CUDA) {
IdentifierInfo *II = NewFD->getIdentifier();
- if (II && II->isStr(getCudaConfigureFuncName()) &&
+ if (II && II->isStr(CUDA().getCudaConfigureFuncName()) &&
!NewFD->isInvalidDecl() &&
NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) {
if (!R->castAs<FunctionType>()->getReturnType()->isScalarType())
Diag(NewFD->getLocation(), diag::err_config_scalar_return)
- << getCudaConfigureFuncName();
+ << CUDA().getCudaConfigureFuncName();
Context.setcudaConfigureCallDecl(NewFD);
}
@@ -12398,7 +12399,7 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
}
if (!Redeclaration && LangOpts.CUDA)
- checkCUDATargetOverload(NewFD, Previous);
+ CUDA().checkCUDATargetOverload(NewFD, Previous);
}
// Check if the function definition uses any AArch64 SME features without
@@ -14415,7 +14416,7 @@ StmtResult Sema::ActOnCXXForRangeIdentifier(Scope *S, SourceLocation IdentLoc,
void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
if (var->isInvalidDecl()) return;
- MaybeAddCUDAConstantAttr(var);
+ CUDA().MaybeAddCUDAConstantAttr(var);
if (getLangOpts().OpenCL) {
// OpenCL v2.0 s6.12.5 - Every block variable declaration must have an
@@ -14829,7 +14830,7 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) {
// variables whether they are local or not. CUDA also allows
// constant initializers for __constant__ and __device__ variables.
if (getLangOpts().CUDA)
- checkAllowedCUDAInitializer(VD);
+ CUDA().checkAllowedCUDAInitializer(VD);
// Grab the dllimport or dllexport attribute off of the VarDecl.
const InheritableAttr *DLLAttr = getDLLAttr(VD);
@@ -20666,7 +20667,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
// when compiling for host, device and global functions are never emitted.
// (Technically, we do emit a host-side stub for global functions, but this
// doesn't count for our purposes here.)
- CUDAFunctionTarget T = IdentifyCUDATarget(FD);
+ CUDAFunctionTarget T = CUDA().IdentifyCUDATarget(FD);
if (LangOpts.CUDAIsDevice && T == CUDAFunctionTarget::Host)
return FunctionEmissionStatus::CUDADiscarded;
if (!LangOpts.CUDAIsDevice &&
@@ -20691,5 +20692,5 @@ bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {
// for host, only HD functions actually called from the host get marked as
// known-emitted.
return LangOpts.CUDA && !LangOpts.CUDAIsDevice &&
- IdentifyCUDATarget(Callee) == CUDAFunctionTarget::Global;
+ CUDA().IdentifyCUDATarget(Callee) == CUDAFunctionTarget::Global;
}
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 56c9d90c9b52b3..0e7e04bca3cf0a 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -39,6 +39,7 @@
#include "clang/Sema/ParsedAttr.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaHLSL.h"
#include "clang/Sema/SemaInternal.h"
#include "llvm/ADT/STLExtras.h"
@@ -5099,8 +5100,8 @@ static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
return;
}
if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
- S.CUDADiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared)
- << llvm::to_underlying(S.CurrentCUDATarget()))
+ S.CUDA().CUDADiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared)
+ << llvm::to_underlying(S.CUDA().CurrentCUDATarget()))
return;
D->addAttr(::new (S.Context) CUDASharedAttr(S.Context, AL));
}
@@ -5190,7 +5191,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
// in the Decl node for syntactic reasoning, e.g., pretty-printing.
CallingConv CC;
if (S.CheckCallingConvAttr(AL, CC, /*FD*/ nullptr,
- S.IdentifyCUDATarget(dyn_cast<FunctionDecl>(D))))
+ S.CUDA().IdentifyCUDATarget(dyn_cast<FunctionDecl>(D))))
return;
if (!isa<ObjCMethodDecl>(D)) {
@@ -5495,7 +5496,7 @@ bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
if (LangOpts.CUDA) {
auto *Aux = Context.getAuxTargetInfo();
assert(FD || CFT != CUDAFunctionTarget::InvalidTarget);
- auto CudaTarget = FD ? IdentifyCUDATarget(FD) : CFT;
+ auto CudaTarget = FD ? CUDA().IdentifyCUDATarget(FD) : CFT;
bool CheckHost = false, CheckDevice = false;
switch (CudaTarget) {
case CUDAFunctionTarget::HostDevice:
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index 51c14443d2d8f1..a1e3f639c7d50c 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -42,6 +42,7 @@
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "llvm/ADT/ArrayRef.h"
@@ -9883,7 +9884,7 @@ bool Sema::ShouldDeleteSpecialMember(CXXMethodDecl *MD,
if (ICI)
RealCSM = getSpecialMember(MD);
- return inferCUDATargetForImplicitSpecialMember(RD, RealCSM, MD,
+ return CUDA().inferCUDATargetForImplicitSpecialMember(RD, RealCSM, MD,
SMI.ConstArg, Diagnose);
}
@@ -14055,7 +14056,7 @@ CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
setupImplicitSpecialMemberType(DefaultCon, Context.VoidTy, std::nullopt);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::DefaultConstructor, DefaultCon,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -14341,7 +14342,7 @@ CXXDestructorDecl *Sema::DeclareImplicitDestructor(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(Destructor, Context.VoidTy, std::nullopt);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::Destructor, Destructor,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -14983,7 +14984,7 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(CopyAssignment, RetType, ArgType);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::CopyAssignment, CopyAssignment,
/* ConstRHS */ Const,
/* Diagnose */ false);
@@ -15335,7 +15336,7 @@ CXXMethodDecl *Sema::DeclareImplicitMoveAssignment(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(MoveAssignment, RetType, ArgType);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::MoveAssignment, MoveAssignment,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -15733,7 +15734,7 @@ CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
setupImplicitSpecialMemberType(CopyConstructor, Context.VoidTy, ArgType);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::CopyConstructor, CopyConstructor,
/* ConstRHS */ Const,
/* Diagnose */ false);
@@ -15878,7 +15879,7 @@ CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
setupImplicitSpecialMemberType(MoveConstructor, Context.VoidTy, ArgType);
if (getLangOpts().CUDA)
- inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferCUDATargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::MoveConstructor, MoveConstructor,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -16184,7 +16185,7 @@ ExprResult Sema::BuildCXXConstructExpr(
DeclInitType->getBaseElementTypeUnsafe()->getAsCXXRecordDecl()) &&
"given constructor for wrong type");
MarkFunctionReferenced(ConstructLoc, Constructor);
- if (getLangOpts().CUDA && !CheckCUDACall(ConstructLoc, Constructor))
+ if (getLangOpts().CUDA && !CUDA().CheckCUDACall(ConstructLoc, Constructor))
return ExprError();
return CheckForImmediateInvocation(
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index b294d2bd9f53f2..63f764843ec0e0 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -49,6 +49,7 @@
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaFixItUtils.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
@@ -308,7 +309,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
DeduceReturnType(FD, Loc))
return true;
- if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
+ if (getLangOpts().CUDA && !CUDA().CheckCUDACall(Loc, FD))
return true;
}
@@ -17307,7 +17308,7 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc,
// CUDA device code does not support varargs.
if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
if (const FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext)) {
- CUDAFunctionTarget T = IdentifyCUDATarget(F);
+ CUDAFunctionTarget T = CUDA().IdentifyCUDATarget(F);
if (T == CUDAFunctionTarget::Global || T == CUDAFunctionTarget::Device ||
T == CUDAFunctionTarget::HostDevice)
return ExprError(Diag(E->getBeginLoc(), diag::err_va_arg_in_device));
@@ -18961,7 +18962,7 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
checkSpecializationReachability(Loc, Func);
if (getLangOpts().CUDA)
- CheckCUDACall(Loc, Func);
+ CUDA().CheckCUDACall(Loc, Func);
// If we need a definition, try to create one.
if (NeedDefinition && !Func->getBody()) {
@@ -19108,7 +19109,7 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
// side. Therefore keep trying until it is recorded.
if (LangOpts.OffloadImplicitHostDeviceTemplates && LangOpts.CUDAIsDevice &&
!getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(Func))
- CUDARecordImplicitHostDeviceFuncUsedByDevice(Func);
+ CUDA().CUDARecordImplicitHostDeviceFuncUsedByDevice(Func);
// If this is the first "real" use, act on that.
if (OdrUse == OdrUseContext::Used && !Func->isUsed(/*CheckUsedAttr=*/false)) {
@@ -19181,9 +19182,9 @@ MarkVarDeclODRUsed(ValueDecl *V, SourceLocation Loc, Sema &SemaRef,
if (SemaRef.LangOpts.CUDA && Var->hasGlobalStorage()) {
auto *FD = dyn_cast_or_null<FunctionDecl>(SemaRef.CurContext);
- auto VarTarget = SemaRef.IdentifyCUDATarget(Var);
- auto UserTarget = SemaRef.IdentifyCUDATarget(FD);
- if (VarTarget == Sema::CVT_Host &&
+ auto VarTarget = SemaRef.CUDA().IdentifyCUDATarget(Var);
+ auto UserTarget = SemaRef.CUDA().IdentifyCUDATarget(FD);
+ if (VarTarget == SemaCUDA::CVT_Host &&
(UserTarget == CUDAFunctionTarget::Device ||
UserTarget == CUDAFunctionTarget::HostDevice ||
UserTarget == CUDAFunctionTarget::Global)) {
@@ -19199,7 +19200,7 @@ MarkVarDeclODRUsed(ValueDecl *V, SourceLocation Loc, Sema &SemaRef,
? diag::note_cuda_const_var_unpromoted
: diag::note_cuda_host_var);
}
- } else if (VarTarget == Sema::CVT_Device &&
+ } else if (VarTarget == SemaCUDA::CVT_Device &&
!Var->hasAttr<CUDASharedAttr>() &&
(UserTarget == CUDAFunctionTarget::Host ||
UserTarget == CUDAFunctionTarget::HostDevice)) {
diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index ce9d5c26e21858..1c80c7467ee1da 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -38,6 +38,7 @@
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/SemaLambda.h"
#include "clang/Sema/Template.h"
@@ -884,8 +885,8 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex,
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
- CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
- << "throw" << llvm::to_underlying(CurrentCUDATarget());
+ CUDA().CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
+ << "throw" << llvm::to_underlying(CUDA().CurrentCUDATarget());
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw";
@@ -1708,17 +1709,17 @@ bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
// [CUDA] Ignore this function, if we can't call it.
const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
if (getLangOpts().CUDA) {
- auto CallPreference = IdentifyCUDAPreference(Caller, Method);
+ auto CallPreference = CUDA().IdentifyCUDAPreference(Caller, Method);
// If it's not callable at all, it's not the right function.
- if (CallPreference < CFP_WrongSide)
+ if (CallPreference < SemaCUDA::CFP_WrongSide)
return false;
- if (CallPreference == CFP_WrongSide) {
+ if (CallPreference == SemaCUDA::CFP_WrongSide) {
// Maybe. We have to check if there are better alternatives.
DeclContext::lookup_result R =
Method->getDeclContext()->lookup(Method->getDeclName());
for (const auto *D : R) {
if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
- if (IdentifyCUDAPreference(Caller, FD) > CFP_WrongSide)
+ if (CUDA().IdentifyCUDAPreference(Caller, FD) > SemaCUDA::CFP_WrongSide)
return false;
}
}
@@ -1737,7 +1738,7 @@ bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
return llvm::none_of(PreventedBy, [&](const FunctionDecl *FD) {
assert(FD->getNumParams() == 1 &&
"Only single-operand functions should be in PreventedBy");
- return IdentifyCUDAPreference(Caller, FD) >= CFP_HostDevice;
+ return CUDA().IdentifyCUDAPreference(Caller, FD) >= SemaCUDA::CFP_HostDevice;
});
}
@@ -1774,7 +1775,7 @@ namespace {
UsualDeallocFnInfo(Sema &S, DeclAccessPair Found)
: Found(Found), FD(dyn_cast<FunctionDecl>(Found->getUnderlyingDecl())),
Destroying(false), HasSizeT(false), HasAlignValT(false),
- CUDAPref(Sema::CFP_Native) {
+ CUDAPref(SemaCUDA::CFP_Native) {
// A function template declaration is never a usual deallocation function.
if (!FD)
return;
@@ -1800,7 +1801,7 @@ namespace {
// In CUDA, determine how much we'd like / dislike to call this.
if (S.getLangOpts().CUDA)
- CUDAPref = S.IdentifyCUDAPreference(
+ CUDAPref = S.CUDA().IdentifyCUDAPreference(
S.getCurFunctionDecl(/*AllowLambda=*/true), FD);
}
@@ -1831,7 +1832,7 @@ namespace {
DeclAccessPair Found;
FunctionDecl *FD;
bool Destroying, HasSizeT, HasAlignValT;
- Sema::CUDAFunctionPreference CUDAPref;
+ SemaCUDA::CUDAFunctionPreference CUDAPref;
};
}
@@ -1855,7 +1856,7 @@ static UsualDeallocFnInfo resolveDeallocationOverload(
for (auto I = R.begin(), E = R.end(); I != E; ++I) {
UsualDeallocFnInfo Info(S, I.getPair());
if (!Info || !isNonPlacementDeallocationFunction(S, Info.FD) ||
- Info.CUDAPref == Sema::CFP_Never)
+ Info.CUDAPref == SemaCUDA::CFP_Never)
continue;
if (!Best) {
@@ -2956,7 +2957,7 @@ bool Sema::FindAllocationFunctions(SourceLocation StartLoc, SourceRange Range,
}
if (getLangOpts().CUDA)
- EraseUnwantedCUDAMatches(getCurFunctionDecl(/*AllowLambda=*/true),
+ CUDA().EraseUnwantedCUDAMatches(getCurFunctionDecl(/*AllowLambda=*/true),
Matches);
} else {
// C++1y [expr.new]p22:
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index 5b95bae567b721..32620675711997 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -18,6 +18,7 @@
#include "clang/Sema/Lookup.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/SemaLambda.h"
#include "clang/Sema/Template.h"
@@ -1393,7 +1394,7 @@ void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
// CUDA lambdas get implicit host and device attributes.
if (getLangOpts().CUDA)
- CUDASetLambdaAttrs(Method);
+ CUDA().CUDASetLambdaAttrs(Method);
// OpenMP lambdas might get assumumption attributes.
if (LangOpts.OpenMP)
@@ -2136,7 +2137,7 @@ ExprResult Sema::BuildLambdaExpr(SourceLocation StartLoc, SourceLocation EndLoc,
CaptureInits.push_back(Init.get());
if (LangOpts.CUDA)
- CUDACheckLambdaCapture(CallOperator, From);
+ CUDA().CUDACheckLambdaCapture(CallOperator, From);
}
Class->setCaptures(Context, Captures);
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index e1155dc2d5d285..1b8b3eedfe18fe 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -31,6 +31,7 @@
#include "clang/Sema/Initialization.h"
#include "clang/Sema/Lookup.h"
#include "clang/Sema/Overload.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "clang/Sema/TemplateDeduction.h"
@@ -1549,8 +1550,8 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New,
// Don't allow overloading of destructors. (In theory we could, but it
// would be a giant change to clang.)
if (!isa<CXXDestructorDecl>(New)) {
- CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(New),
- OldTarget = SemaRef.IdentifyCUDATarget(Old);
+ CUDAFunctionTarget NewTarget = SemaRef.CUDA().IdentifyCUDATarget(New),
+ OldTarget = SemaRef.CUDA().IdentifyCUDATarget(Old);
if (NewTarget != CUDAFunctionTarget::InvalidTarget) {
assert((OldTarget != CUDAFunctionTarget::InvalidTarget) &&
"Unexpected invalid target.");
@@ -7100,7 +7101,7 @@ void Sema::AddOverloadCandidate(
// inferred for the member automatically, based on the bases and fields of
// the class.
if (!(Caller && Caller->isImplicit()) &&
- !IsAllowedCUDACall(Caller, Function)) {
+ !CUDA().IsAllowedCUDACall(Caller, Function)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -7618,7 +7619,7 @@ Sema::AddMethodCandidate(CXXMethodDecl *Method, DeclAccessPair FoundDecl,
// (CUDA B.1): Check for invalid calls between targets.
if (getLangOpts().CUDA)
- if (!IsAllowedCUDACall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
+ if (!CUDA().IsAllowedCUDACall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -10452,14 +10453,14 @@ bool clang::isBetterOverloadCandidate(
if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function &&
S.getLangOpts().GPUExcludeWrongSideOverloads) {
if (FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true)) {
- bool IsCallerImplicitHD = Sema::isCUDAImplicitHostDeviceFunction(Caller);
+ bool IsCallerImplicitHD = SemaCUDA::isCUDAImplicitHostDeviceFunction(Caller);
bool IsCand1ImplicitHD =
- Sema::isCUDAImplicitHostDeviceFunction(Cand1.Function);
+ SemaCUDA::isCUDAImplicitHostDeviceFunction(Cand1.Function);
bool IsCand2ImplicitHD =
- Sema::isCUDAImplicitHostDeviceFunction(Cand2.Function);
- auto P1 = S.IdentifyCUDAPreference(Caller, Cand1.Function);
- auto P2 = S.IdentifyCUDAPreference(Caller, Cand2.Function);
- assert(P1 != Sema::CFP_Never && P2 != Sema::CFP_Never);
+ SemaCUDA::isCUDAImplicitHostDeviceFunction(Cand2.Function);
+ auto P1 = S.CUDA().IdentifyCUDAPreference(Caller, Cand1.Function);
+ auto P2 = S.CUDA().IdentifyCUDAPreference(Caller, Cand2.Function);
+ assert(P1 != SemaCUDA::CFP_Never && P2 != SemaCUDA::CFP_Never);
// The implicit HD function may be a function in a system header which
// is forced by pragma. In device compilation, if we prefer HD candidates
// over wrong-sided candidates, overloading resolution may change, which
@@ -10473,8 +10474,8 @@ bool clang::isBetterOverloadCandidate(
auto EmitThreshold =
(S.getLangOpts().CUDAIsDevice && IsCallerImplicitHD &&
(IsCand1ImplicitHD || IsCand2ImplicitHD))
- ? Sema::CFP_Never
- : Sema::CFP_WrongSide;
+ ? SemaCUDA::CFP_Never
+ : SemaCUDA::CFP_WrongSide;
auto Cand1Emittable = P1 > EmitThreshold;
auto Cand2Emittable = P2 > EmitThreshold;
if (Cand1Emittable && !Cand2Emittable)
@@ -10758,8 +10759,8 @@ bool clang::isBetterOverloadCandidate(
// to determine which is better.
if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
- return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
- S.IdentifyCUDAPreference(Caller, Cand2.Function);
+ return S.CUDA().IdentifyCUDAPreference(Caller, Cand1.Function) >
+ S.CUDA().IdentifyCUDAPreference(Caller, Cand2.Function);
}
// General member function overloading is handled above, so this only handles
@@ -10891,15 +10892,15 @@ OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
// Check viable function only.
return Cand->Viable && Cand->Function &&
- S.IdentifyCUDAPreference(Caller, Cand->Function) ==
- Sema::CFP_SameSide;
+ S.CUDA().IdentifyCUDAPreference(Caller, Cand->Function) ==
+ SemaCUDA::CFP_SameSide;
});
if (ContainsSameSideCandidate) {
auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {
// Check viable function only to avoid unnecessary data copying/moving.
return Cand->Viable && Cand->Function &&
- S.IdentifyCUDAPreference(Caller, Cand->Function) ==
- Sema::CFP_WrongSide;
+ S.CUDA().IdentifyCUDAPreference(Caller, Cand->Function) ==
+ SemaCUDA::CFP_WrongSide;
};
llvm::erase_if(Candidates, IsWrongSideCandidate);
}
@@ -11938,8 +11939,8 @@ static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
FunctionDecl *Callee = Cand->Function;
- CUDAFunctionTarget CallerTarget = S.IdentifyCUDATarget(Caller),
- CalleeTarget = S.IdentifyCUDATarget(Callee);
+ CUDAFunctionTarget CallerTarget = S.CUDA().IdentifyCUDATarget(Caller),
+ CalleeTarget = S.CUDA().IdentifyCUDATarget(Callee);
std::string FnDesc;
std::pair<OverloadCandidateKind, OverloadCandidateSelect> FnKindPair =
@@ -11986,7 +11987,7 @@ static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
}
}
- S.inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth,
+ S.CUDA().inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth,
/* ConstRHS */ ConstRHS,
/* Diagnose */ true);
}
@@ -13060,7 +13061,7 @@ class AddressOfFunctionResolver {
if (S.getLangOpts().CUDA) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
if (!(Caller && Caller->isImplicit()) &&
- !S.IsAllowedCUDACall(Caller, FunDecl))
+ !S.CUDA().IsAllowedCUDACall(Caller, FunDecl))
return false;
}
if (FunDecl->isMultiVersion()) {
@@ -13180,7 +13181,7 @@ class AddressOfFunctionResolver {
}
void EliminateSuboptimalCudaMatches() {
- S.EraseUnwantedCUDAMatches(S.getCurFunctionDecl(/*AllowLambda=*/true),
+ S.CUDA().EraseUnwantedCUDAMatches(S.getCurFunctionDecl(/*AllowLambda=*/true),
Matches);
}
@@ -13335,8 +13336,8 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
// Return positive for better, negative for worse, 0 for equal preference.
auto CheckCUDAPreference = [&](FunctionDecl *FD1, FunctionDecl *FD2) {
FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
- return static_cast<int>(IdentifyCUDAPreference(Caller, FD1)) -
- static_cast<int>(IdentifyCUDAPreference(Caller, FD2));
+ return static_cast<int>(CUDA().IdentifyCUDAPreference(Caller, FD1)) -
+ static_cast<int>(CUDA().IdentifyCUDAPreference(Caller, FD2));
};
auto CheckMoreConstrained = [&](FunctionDecl *FD1,
diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 1c2f6120f6218b..732e34257c47d8 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -33,6 +33,7 @@
#include "clang/Sema/Ownership.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseMap.h"
@@ -4574,8 +4575,8 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
- CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
- << "try" << llvm::to_underlying(CurrentCUDATarget());
+ CUDA().CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
+ << "try" << llvm::to_underlying(CUDA().CurrentCUDATarget());
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try";
diff --git a/clang/lib/Sema/SemaTemplate.cpp b/clang/lib/Sema/SemaTemplate.cpp
index e0f5e53dc2481e..fa1127b10aa7e2 100644
--- a/clang/lib/Sema/SemaTemplate.cpp
+++ b/clang/lib/Sema/SemaTemplate.cpp
@@ -33,6 +33,7 @@
#include "clang/Sema/Overload.h"
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "clang/Sema/TemplateDeduction.h"
@@ -10155,9 +10156,9 @@ bool Sema::CheckFunctionTemplateSpecialization(
// take target attributes into account, we reject candidates
// here that have a different target.
if (LangOpts.CUDA &&
- IdentifyCUDATarget(Specialization,
+ CUDA().IdentifyCUDATarget(Specialization,
/* IgnoreImplicitHDAttr = */ true) !=
- IdentifyCUDATarget(FD, /* IgnoreImplicitHDAttr = */ true)) {
+ CUDA().IdentifyCUDATarget(FD, /* IgnoreImplicitHDAttr = */ true)) {
FailedCandidates.addCandidate().set(
I.getPair(), FunTmpl->getTemplatedDecl(),
MakeDeductionFailureInfo(
@@ -10328,7 +10329,7 @@ bool Sema::CheckFunctionTemplateSpecialization(
// virtue e.g. of being constexpr, and it passes these implicit
// attributes on to its specializations.)
if (LangOpts.CUDA)
- inheritCUDATargetAttrs(FD, *Specialization->getPrimaryTemplate());
+ CUDA().inheritCUDATargetAttrs(FD, *Specialization->getPrimaryTemplate());
// The "previous declaration" for this function template specialization is
// the prior function template specialization.
@@ -11364,9 +11365,9 @@ DeclResult Sema::ActOnExplicitInstantiation(Scope *S,
// target attributes into account, we reject candidates here that
// have a different target.
if (LangOpts.CUDA &&
- IdentifyCUDATarget(Specialization,
+ CUDA().IdentifyCUDATarget(Specialization,
/* IgnoreImplicitHDAttr = */ true) !=
- IdentifyCUDATarget(D.getDeclSpec().getAttributes())) {
+ CUDA().IdentifyCUDATarget(D.getDeclSpec().getAttributes())) {
FailedCandidates.addCandidate().set(
P.getPair(), FunTmpl->getTemplatedDecl(),
MakeDeductionFailureInfo(
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 7f510d34d671ee..15e247a956e3ed 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -26,6 +26,7 @@
#include "clang/Sema/Initialization.h"
#include "clang/Sema/Lookup.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "clang/Sema/TemplateInstCallback.h"
@@ -5537,7 +5538,7 @@ void Sema::InstantiateVariableInitializer(
}
if (getLangOpts().CUDA)
- checkAllowedCUDAInitializer(Var);
+ CUDA().checkAllowedCUDAInitializer(Var);
}
/// Instantiate the definition of the given variable from its
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index b3f6078952f6eb..09769ced5d9c21 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -33,6 +33,7 @@
#include "clang/Sema/Lookup.h"
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/ScopeInfo.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "clang/Sema/TemplateInstCallback.h"
@@ -2735,7 +2736,7 @@ QualType Sema::BuildArrayType(QualType T, ArraySizeModifier ASM,
bool IsCUDADevice = (getLangOpts().CUDA && getLangOpts().CUDAIsDevice);
targetDiag(Loc,
IsCUDADevice ? diag::err_cuda_vla : diag::err_vla_unsupported)
- << (IsCUDADevice ? llvm::to_underlying(CurrentCUDATarget()) : 0);
+ << (IsCUDADevice ? llvm::to_underlying(CUDA().CurrentCUDATarget()) : 0);
} else if (sema::FunctionScopeInfo *FSI = getCurFunction()) {
// VLAs are supported on this target, but we may need to do delayed
// checking that the VLA is not being used within a coroutine.
@@ -3618,7 +3619,7 @@ static QualType GetDeclSpecTypeForDeclarator(TypeProcessingState &state,
// D.getDeclarationAttributes()) because those are always C++11 attributes,
// and those don't get distributed.
distributeTypeAttrsFromDeclarator(
- state, T, SemaRef.IdentifyCUDATarget(D.getAttributes()));
+ state, T, SemaRef.CUDA().IdentifyCUDATarget(D.getAttributes()));
// Find the deduced type in this type. Look in the trailing return type if we
// have one, otherwise in the DeclSpec type.
@@ -4139,7 +4140,7 @@ static CallingConv getCCForDeclaratorChunk(
// handleFunctionTypeAttr.
CallingConv CC;
if (!S.CheckCallingConvAttr(AL, CC, /*FunctionDecl=*/nullptr,
- S.IdentifyCUDATarget(D.getAttributes())) &&
+ S.CUDA().IdentifyCUDATarget(D.getAttributes())) &&
(!FTI.isVariadic || supportsVariadicCall(CC))) {
return CC;
}
@@ -5825,7 +5826,7 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
// See if there are any attributes on this declarator chunk.
processTypeAttrs(state, T, TAL_DeclChunk, DeclType.getAttrs(),
- S.IdentifyCUDATarget(D.getAttributes()));
+ S.CUDA().IdentifyCUDATarget(D.getAttributes()));
if (DeclType.Kind != DeclaratorChunk::Paren) {
if (ExpectNoDerefChunk && !IsNoDerefableChunk(DeclType))
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index ce96ce2bdbcce6..93051d05cc4e8c 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -78,6 +78,7 @@
#include "clang/Sema/ObjCMethodList.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/Sema.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/Weak.h"
#include "clang/Serialization/ASTBitCodes.h"
#include "clang/Serialization/ASTDeserializationListener.h"
@@ -8274,7 +8275,7 @@ void ASTReader::UpdateSema() {
PragmaMSPointersToMembersState,
PointersToMembersPragmaLocation);
}
- SemaObj->ForceCUDAHostDeviceDepth = ForceCUDAHostDeviceDepth;
+ SemaObj->CUDA().ForceCUDAHostDeviceDepth = ForceCUDAHostDeviceDepth;
if (PragmaAlignPackCurrentValue) {
// The bottom of the stack might have a default value. It must be adjusted
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index d9ba10ab608783..d959bb5e69d03b 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -65,6 +65,7 @@
#include "clang/Sema/IdentifierResolver.h"
#include "clang/Sema/ObjCMethodList.h"
#include "clang/Sema/Sema.h"
+#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/Weak.h"
#include "clang/Serialization/ASTBitCodes.h"
#include "clang/Serialization/ASTReader.h"
@@ -4335,8 +4336,8 @@ void ASTWriter::WriteOpenCLExtensions(Sema &SemaRef) {
Stream.EmitRecord(OPENCL_EXTENSIONS, Record);
}
void ASTWriter::WriteCUDAPragmas(Sema &SemaRef) {
- if (SemaRef.ForceCUDAHostDeviceDepth > 0) {
- RecordData::value_type Record[] = {SemaRef.ForceCUDAHostDeviceDepth};
+ if (SemaRef.CUDA().ForceCUDAHostDeviceDepth > 0) {
+ RecordData::value_type Record[] = {SemaRef.CUDA().ForceCUDAHostDeviceDepth};
Stream.EmitRecord(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH, Record);
}
}
>From a564a25f423e62ba0e0f8f77a65c57e22114ca27 Mon Sep 17 00:00:00 2001
From: Vlad Serebrennikov <serebrennikov.vladislav at gmail.com>
Date: Fri, 12 Apr 2024 21:41:32 +0300
Subject: [PATCH 2/4] Drop `CUDA` from names of `SemaCUDA` functions
---
clang/include/clang/Sema/SemaBase.h | 2 +-
clang/include/clang/Sema/SemaCUDA.h | 64 +++++-----
clang/include/clang/Serialization/ASTReader.h | 2 +-
clang/lib/Parse/ParseExpr.cpp | 2 +-
clang/lib/Parse/ParsePragma.cpp | 4 +-
clang/lib/Sema/Sema.cpp | 8 +-
clang/lib/Sema/SemaBase.cpp | 4 +-
clang/lib/Sema/SemaCUDA.cpp | 112 +++++++++---------
clang/lib/Sema/SemaDecl.cpp | 18 +--
clang/lib/Sema/SemaDeclAttr.cpp | 8 +-
clang/lib/Sema/SemaDeclCXX.cpp | 18 +--
clang/lib/Sema/SemaExpr.cpp | 12 +-
clang/lib/Sema/SemaExprCXX.cpp | 14 +--
clang/lib/Sema/SemaLambda.cpp | 4 +-
clang/lib/Sema/SemaOverload.cpp | 42 +++----
clang/lib/Sema/SemaStmt.cpp | 4 +-
clang/lib/Sema/SemaTemplate.cpp | 10 +-
.../lib/Sema/SemaTemplateInstantiateDecl.cpp | 2 +-
clang/lib/Sema/SemaType.cpp | 8 +-
clang/lib/Serialization/ASTReader.cpp | 4 +-
clang/lib/Serialization/ASTWriter.cpp | 4 +-
21 files changed, 173 insertions(+), 173 deletions(-)
diff --git a/clang/include/clang/Sema/SemaBase.h b/clang/include/clang/Sema/SemaBase.h
index ff718022fca03c..3220f71dd797ed 100644
--- a/clang/include/clang/Sema/SemaBase.h
+++ b/clang/include/clang/Sema/SemaBase.h
@@ -146,7 +146,7 @@ class SemaBase {
/// if (SemaDiagnosticBuilder(...) << foo << bar)
/// return ExprError();
///
- /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+ /// But see DiagIfDeviceCode() and DiagIfHostCode() -- you probably
/// want to use these instead of creating a SemaDiagnosticBuilder yourself.
operator bool() const { return isImmediate(); }
diff --git a/clang/include/clang/Sema/SemaCUDA.h b/clang/include/clang/Sema/SemaCUDA.h
index 1ceaab06a5cc62..b7dd57aaeadcd0 100644
--- a/clang/include/clang/Sema/SemaCUDA.h
+++ b/clang/include/clang/Sema/SemaCUDA.h
@@ -39,14 +39,14 @@ class SemaCUDA : public SemaBase {
/// 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 PushForceCUDAHostDevice();
+ 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 PopForceCUDAHostDevice();
+ bool PopForceHostDevice();
- ExprResult ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
+ ExprResult ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc);
@@ -57,7 +57,7 @@ class SemaCUDA : public SemaBase {
SourceLocation Loc;
};
- /// FunctionDecls and SourceLocations for which CheckCUDACall has emitted a
+ /// 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;
@@ -85,26 +85,26 @@ class SemaCUDA : public SemaBase {
/// Example usage:
///
/// // Variable-length arrays are not allowed in CUDA device code.
- /// if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
+ /// if (DiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentTarget())
/// return ExprError();
/// // Otherwise, continue parsing as normal.
- SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
+ SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc,
unsigned DiagID);
/// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
/// context is "used as host code".
///
- /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
- SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+ /// 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 IdentifyCUDATarget(const FunctionDecl *D,
+ CUDAFunctionTarget IdentifyTarget(const FunctionDecl *D,
bool IgnoreImplicitHDAttr = false);
- CUDAFunctionTarget IdentifyCUDATarget(const ParsedAttributesView &Attrs);
+ CUDAFunctionTarget IdentifyTarget(const ParsedAttributesView &Attrs);
enum CUDAVariableTarget {
CVT_Device, /// Emitted on device side with a shadow variable on host side
@@ -113,7 +113,7 @@ class SemaCUDA : public SemaBase {
CVT_Unified, /// Emitted as a unified address, e.g. managed variables
};
/// Determines whether the given variable is emitted on host or device side.
- CUDAVariableTarget IdentifyCUDATarget(const VarDecl *D);
+ CUDAVariableTarget IdentifyTarget(const VarDecl *D);
/// Defines kinds of CUDA global host/device context where a function may be
/// called.
@@ -139,11 +139,11 @@ class SemaCUDA : public SemaBase {
};
/// Gets the CUDA target for the current context.
- CUDAFunctionTarget CurrentCUDATarget() {
- return IdentifyCUDATarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
+ CUDAFunctionTarget CurrentTarget() {
+ return IdentifyTarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
}
- static bool isCUDAImplicitHostDeviceFunction(const FunctionDecl *D);
+ static bool isImplicitHostDeviceFunction(const FunctionDecl *D);
// CUDA function call preference. Must be ordered numerically from
// worst to best.
@@ -165,7 +165,7 @@ class SemaCUDA : public SemaBase {
/// \param Callee target function
///
/// \returns preference value for particular Caller/Callee combination.
- CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
+ CUDAFunctionPreference IdentifyPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee);
/// Determines whether Caller may invoke Callee, based on their CUDA
@@ -173,19 +173,19 @@ class SemaCUDA : public SemaBase {
///
/// 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 IsAllowedCUDACall(const FunctionDecl *Caller,
+ bool IsAllowedCall(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
- return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
+ return IdentifyPreference(Caller, Callee) != CFP_Never;
}
/// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
/// depending on FD and the current compilation settings.
- void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
+ void maybeAddHostDeviceAttrs(FunctionDecl *FD,
const LookupResult &Previous);
/// May add implicit CUDAConstantAttr attribute to VD, depending on VD
/// and current compilation settings.
- void MaybeAddCUDAConstantAttr(VarDecl *VD);
+ void MaybeAddConstantAttr(VarDecl *VD);
/// Check whether we're allowed to call Callee from the current context.
///
@@ -201,25 +201,25 @@ class SemaCUDA : public SemaBase {
/// deferred errors.
///
/// - Otherwise, returns true without emitting any diagnostics.
- bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
+ bool CheckCall(SourceLocation Loc, FunctionDecl *Callee);
- void CUDACheckLambdaCapture(CXXMethodDecl *D, const sema::Capture &Capture);
+ 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 CUDASetLambdaAttrs(CXXMethodDecl *Method);
+ 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 CUDARecordImplicitHostDeviceFuncUsedByDevice(const FunctionDecl *FD);
+ 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 EraseUnwantedCUDAMatches(
+ void EraseUnwantedMatches(
const FunctionDecl *Caller,
llvm::SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches);
@@ -234,15 +234,15 @@ class SemaCUDA : public SemaBase {
/// \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 inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
+ 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 isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
- bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
+ 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.
@@ -252,22 +252,22 @@ class SemaCUDA : public SemaBase {
// __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 checkAllowedCUDAInitializer(VarDecl *VD);
+ void checkAllowedInitializer(VarDecl *VD);
/// Check whether NewFD is a valid overload for CUDA. Emits
/// diagnostics and invalidates NewFD if not.
- void checkCUDATargetOverload(FunctionDecl *NewFD,
+ void checkTargetOverload(FunctionDecl *NewFD,
const LookupResult &Previous);
/// Copies target attributes from the template TD to the function FD.
- void inheritCUDATargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD);
+ 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 getCudaConfigureFuncName() const;
+ std::string getConfigureFuncName() const;
private:
- unsigned ForceCUDAHostDeviceDepth = 0;
+ unsigned ForceHostDeviceDepth = 0;
friend class ASTReader;
friend class ASTWriter;
diff --git a/clang/include/clang/Serialization/ASTReader.h b/clang/include/clang/Serialization/ASTReader.h
index 6656c1c58dec9d..e3fde887f99cb7 100644
--- a/clang/include/clang/Serialization/ASTReader.h
+++ b/clang/include/clang/Serialization/ASTReader.h
@@ -873,7 +873,7 @@ class ASTReader
/// Our current depth in #pragma cuda force_host_device begin/end
/// macros.
- unsigned ForceCUDAHostDeviceDepth = 0;
+ unsigned ForceHostDeviceDepth = 0;
/// The IDs of the declarations Sema stores directly.
///
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index e2410555f5e0de..ec05ce616ca647 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -2130,7 +2130,7 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
}
if (!LHS.isInvalid()) {
- ExprResult ECResult = Actions.CUDA().ActOnCUDAExecConfigExpr(getCurScope(),
+ ExprResult ECResult = Actions.CUDA().ActOnExecConfigExpr(getCurScope(),
OpenLoc,
ExecConfigExprs,
CloseLoc);
diff --git a/clang/lib/Parse/ParsePragma.cpp b/clang/lib/Parse/ParsePragma.cpp
index cd784eb0fb0cdc..3979f75b6020db 100644
--- a/clang/lib/Parse/ParsePragma.cpp
+++ b/clang/lib/Parse/ParsePragma.cpp
@@ -3901,8 +3901,8 @@ void PragmaForceCUDAHostDeviceHandler::HandlePragma(
}
if (Info->isStr("begin"))
- Actions.CUDA().PushForceCUDAHostDevice();
- else if (!Actions.CUDA().PopForceCUDAHostDevice())
+ Actions.CUDA().PushForceHostDevice();
+ else if (!Actions.CUDA().PopForceHostDevice())
PP.Diag(FirstTok.getLocation(),
diag::err_pragma_cannot_end_force_cuda_host_device);
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 50de4f33ba8b69..8de202f4f7a0c3 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1838,8 +1838,8 @@ void Sema::emitDeferredDiags() {
// which other not-known-emitted functions.
//
// When we see something which is illegal if the current function is emitted
-// (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or
-// CheckCUDACall), we first check if the current function is known-emitted. If
+// (usually by way of DiagIfDeviceCode, DiagIfHostCode, or
+// CheckCall), we first check if the current function is known-emitted. If
// so, we immediately output the diagnostic.
//
// Otherwise, we "defer" the diagnostic. It sits in Sema::DeviceDeferredDiags
@@ -1902,8 +1902,8 @@ Sema::targetDiag(SourceLocation Loc, unsigned DiagID, const FunctionDecl *FD) {
? diagIfOpenMPDeviceCode(Loc, DiagID, FD)
: diagIfOpenMPHostCode(Loc, DiagID, FD);
if (getLangOpts().CUDA)
- return getLangOpts().CUDAIsDevice ? CUDA().CUDADiagIfDeviceCode(Loc, DiagID)
- : CUDA().CUDADiagIfHostCode(Loc, DiagID);
+ return getLangOpts().CUDAIsDevice ? CUDA().DiagIfDeviceCode(Loc, DiagID)
+ : CUDA().DiagIfHostCode(Loc, DiagID);
if (getLangOpts().SYCLIsDevice)
return SYCL().DiagIfDeviceCode(Loc, DiagID);
diff --git a/clang/lib/Sema/SemaBase.cpp b/clang/lib/Sema/SemaBase.cpp
index 17cff0e84bb0d3..0442fb2929e3c6 100644
--- a/clang/lib/Sema/SemaBase.cpp
+++ b/clang/lib/Sema/SemaBase.cpp
@@ -71,8 +71,8 @@ Sema::SemaDiagnosticBuilder SemaBase::Diag(SourceLocation Loc, unsigned DiagID,
}
SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice
- ? SemaRef.CUDA().CUDADiagIfDeviceCode(Loc, DiagID)
- : SemaRef.CUDA().CUDADiagIfHostCode(Loc, DiagID);
+ ? SemaRef.CUDA().DiagIfDeviceCode(Loc, DiagID)
+ : SemaRef.CUDA().DiagIfHostCode(Loc, DiagID);
SetIsLastErrorImmediate(DB.isImmediate());
return DB;
}
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 6e94e4e8091f60..13d318fe911766 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -38,26 +38,26 @@ template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) {
return false;
}
-void SemaCUDA::PushForceCUDAHostDevice() {
+void SemaCUDA::PushForceHostDevice() {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- ForceCUDAHostDeviceDepth++;
+ ForceHostDeviceDepth++;
}
-bool SemaCUDA::PopForceCUDAHostDevice() {
+bool SemaCUDA::PopForceHostDevice() {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- if (ForceCUDAHostDeviceDepth == 0)
+ if (ForceHostDeviceDepth == 0)
return false;
- ForceCUDAHostDeviceDepth--;
+ ForceHostDeviceDepth--;
return true;
}
-ExprResult SemaCUDA::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
+ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc) {
FunctionDecl *ConfigDecl = getASTContext().getcudaConfigureCallDecl();
if (!ConfigDecl)
return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
- << getCudaConfigureFuncName());
+ << getConfigureFuncName());
QualType ConfigQTy = ConfigDecl->getType();
DeclRefExpr *ConfigDR = new (getASTContext())
@@ -68,7 +68,7 @@ ExprResult SemaCUDA::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
/*IsExecConfig=*/true);
}
-CUDAFunctionTarget SemaCUDA::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
+CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
bool HasHostAttr = false;
bool HasDeviceAttr = false;
bool HasGlobalAttr = false;
@@ -133,8 +133,8 @@ SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(SemaCUDA &S_,
}
}
-/// IdentifyCUDATarget - Determine the CUDA compilation target for this function
-CUDAFunctionTarget SemaCUDA::IdentifyCUDATarget(const FunctionDecl *D,
+/// IdentifyTarget - Determine the CUDA compilation target for this function
+CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
bool IgnoreImplicitHDAttr) {
// Code that lives outside a function gets the target from CurCUDATargetCtx.
if (D == nullptr)
@@ -163,7 +163,7 @@ CUDAFunctionTarget SemaCUDA::IdentifyCUDATarget(const FunctionDecl *D,
}
/// IdentifyTarget - Determine the CUDA compilation target for this variable.
-SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyCUDATarget(const VarDecl *Var) {
+SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyTarget(const VarDecl *Var) {
if (Var->hasAttr<HIPManagedAttr>())
return CVT_Unified;
// Only constexpr and const variabless with implicit constant attribute
@@ -183,7 +183,7 @@ SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyCUDATarget(const VarDecl *Var) {
// - on both sides in host device functions
// - on device side in device or global functions
if (auto *FD = dyn_cast<FunctionDecl>(Var->getDeclContext())) {
- switch (IdentifyCUDATarget(FD)) {
+ switch (IdentifyTarget(FD)) {
case CUDAFunctionTarget::HostDevice:
return CVT_Both;
case CUDAFunctionTarget::Device:
@@ -225,20 +225,20 @@ SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyCUDATarget(const VarDecl *Var) {
// | hd | hd | HD | HD | (b) |
SemaCUDA::CUDAFunctionPreference
-SemaCUDA::IdentifyCUDAPreference(const FunctionDecl *Caller,
+SemaCUDA::IdentifyPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
assert(Callee && "Callee must be valid.");
// Treat ctor/dtor as host device function in device var initializer to allow
// trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
- // will be diagnosed by checkAllowedCUDAInitializer.
+ // will be diagnosed by checkAllowedInitializer.
if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
CurCUDATargetCtx.Target == CUDAFunctionTarget::Device &&
(isa<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee)))
return CFP_HostDevice;
- CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
- CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
+ CUDAFunctionTarget CallerTarget = IdentifyTarget(Caller);
+ CUDAFunctionTarget CalleeTarget = IdentifyTarget(Callee);
// If one of the targets is invalid, the check always fails, no matter what
// the other target is.
@@ -312,13 +312,13 @@ template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) {
return D->isImplicit();
}
-bool SemaCUDA::isCUDAImplicitHostDeviceFunction(const FunctionDecl *D) {
+bool SemaCUDA::isImplicitHostDeviceFunction(const FunctionDecl *D) {
bool IsImplicitDevAttr = hasImplicitAttr<CUDADeviceAttr>(D);
bool IsImplicitHostAttr = hasImplicitAttr<CUDAHostAttr>(D);
return IsImplicitDevAttr && IsImplicitHostAttr;
}
-void SemaCUDA::EraseUnwantedCUDAMatches(
+void SemaCUDA::EraseUnwantedMatches(
const FunctionDecl *Caller,
SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
if (Matches.size() <= 1)
@@ -328,7 +328,7 @@ void SemaCUDA::EraseUnwantedCUDAMatches(
// Gets the CUDA function preference for a call from Caller to Match.
auto GetCFP = [&](const Pair &Match) {
- return IdentifyCUDAPreference(Caller, Match.second);
+ return IdentifyPreference(Caller, Match.second);
};
// Find the best call preference among the functions in Matches.
@@ -370,7 +370,7 @@ resolveCalleeCUDATargetConflict(CUDAFunctionTarget Target1,
return false;
}
-bool SemaCUDA::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
+bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXSpecialMemberKind CSM,
CXXMethodDecl *MemberDecl,
bool ConstRHS,
@@ -425,7 +425,7 @@ bool SemaCUDA::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
if (!SMOR.getMethod())
continue;
- CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR.getMethod());
+ CUDAFunctionTarget BaseMethodTarget = IdentifyTarget(SMOR.getMethod());
if (!InferredTarget) {
InferredTarget = BaseMethodTarget;
} else {
@@ -469,7 +469,7 @@ bool SemaCUDA::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
continue;
CUDAFunctionTarget FieldMethodTarget =
- IdentifyCUDATarget(SMOR.getMethod());
+ IdentifyTarget(SMOR.getMethod());
if (!InferredTarget) {
InferredTarget = FieldMethodTarget;
} else {
@@ -509,7 +509,7 @@ bool SemaCUDA::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
return false;
}
-bool SemaCUDA::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
+bool SemaCUDA::isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
if (!CD->isDefined() && CD->isTemplateInstantiation())
SemaRef.InstantiateFunctionDefinition(Loc, CD->getFirstDecl());
@@ -539,7 +539,7 @@ bool SemaCUDA::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD
if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
if (const CXXConstructExpr *CE =
dyn_cast<CXXConstructExpr>(CI->getInit()))
- return isEmptyCudaConstructor(Loc, CE->getConstructor());
+ return isEmptyConstructor(Loc, CE->getConstructor());
return false;
}))
return false;
@@ -547,7 +547,7 @@ bool SemaCUDA::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD
return true;
}
-bool SemaCUDA::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
+bool SemaCUDA::isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
// No destructor -> no problem.
if (!DD)
return true;
@@ -582,7 +582,7 @@ bool SemaCUDA::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD)
// destructors for all base classes...
if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
- return isEmptyCudaDestructor(Loc, RD->getDestructor());
+ return isEmptyDestructor(Loc, RD->getDestructor());
return true;
}))
return false;
@@ -592,7 +592,7 @@ bool SemaCUDA::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD)
if (CXXRecordDecl *RD = Field->getType()
->getBaseElementTypeUnsafe()
->getAsCXXRecordDecl())
- return isEmptyCudaDestructor(Loc, RD->getDestructor());
+ return isEmptyDestructor(Loc, RD->getDestructor());
return true;
}))
return false;
@@ -632,7 +632,7 @@ bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
if (!Init)
return true;
if (const auto *CE = dyn_cast<CXXConstructExpr>(Init)) {
- return S.isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+ return S.isEmptyConstructor(VD->getLocation(), CE->getConstructor());
}
return false;
};
@@ -645,7 +645,7 @@ bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
};
auto HasEmptyDtor = [&](VarDecl *VD) {
if (const auto *RD = VD->getType()->getAsCXXRecordDecl())
- return S.isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+ return S.isEmptyDestructor(VD->getLocation(), RD->getDestructor());
return true;
};
if (CheckKind == CICK_Shared)
@@ -655,7 +655,7 @@ bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
}
} // namespace
-void SemaCUDA::checkAllowedCUDAInitializer(VarDecl *VD) {
+void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
// Return early if VD is inside a non-instantiated template function since
// the implicit constructor is not defined yet.
if (const FunctionDecl *FD =
@@ -691,7 +691,7 @@ void SemaCUDA::checkAllowedCUDAInitializer(VarDecl *VD) {
InitFn = CE->getDirectCallee();
}
if (InitFn) {
- CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
+ CUDAFunctionTarget InitFnTarget = IdentifyTarget(InitFn);
if (InitFnTarget != CUDAFunctionTarget::Host &&
InitFnTarget != CUDAFunctionTarget::HostDevice) {
Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
@@ -703,22 +703,22 @@ void SemaCUDA::checkAllowedCUDAInitializer(VarDecl *VD) {
}
}
-void SemaCUDA::CUDARecordImplicitHostDeviceFuncUsedByDevice(
+void SemaCUDA::RecordImplicitHostDeviceFuncUsedByDevice(
const FunctionDecl *Callee) {
FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
if (!Caller)
return;
- if (!isCUDAImplicitHostDeviceFunction(Callee))
+ if (!isImplicitHostDeviceFunction(Callee))
return;
- CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
+ CUDAFunctionTarget CallerTarget = IdentifyTarget(Caller);
// Record whether an implicit host device function is used on device side.
if (CallerTarget != CUDAFunctionTarget::Device &&
CallerTarget != CUDAFunctionTarget::Global &&
(CallerTarget != CUDAFunctionTarget::HostDevice ||
- (isCUDAImplicitHostDeviceFunction(Caller) &&
+ (isImplicitHostDeviceFunction(Caller) &&
!getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(Caller))))
return;
@@ -734,14 +734,14 @@ void SemaCUDA::CUDARecordImplicitHostDeviceFuncUsedByDevice(
// system header, in which case we leave the constexpr function unattributed.
//
// In addition, all function decls are treated as __host__ __device__ when
-// ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
+// ForceHostDeviceDepth > 0 (corresponding to code within a
// #pragma clang force_cuda_host_device_begin/end
// pair).
-void SemaCUDA::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
+void SemaCUDA::maybeAddHostDeviceAttrs(FunctionDecl *NewD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- if (ForceCUDAHostDeviceDepth > 0) {
+ if (ForceHostDeviceDepth > 0) {
if (!NewD->hasAttr<CUDAHostAttr>())
NewD->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
if (!NewD->hasAttr<CUDADeviceAttr>())
@@ -801,7 +801,7 @@ void SemaCUDA::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
// TODO: `__constant__` memory may be a limited resource for certain targets.
// A safeguard may be needed at the end of compilation pipeline if
// `__constant__` memory usage goes beyond limit.
-void SemaCUDA::MaybeAddCUDAConstantAttr(VarDecl *VD) {
+void SemaCUDA::MaybeAddConstantAttr(VarDecl *VD) {
// Do not promote dependent variables since the cotr/dtor/initializer are
// not determined. Do it after instantiation.
if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
@@ -815,14 +815,14 @@ void SemaCUDA::MaybeAddCUDAConstantAttr(VarDecl *VD) {
}
}
-SemaBase::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfDeviceCode(SourceLocation Loc,
+SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
- switch (CurrentCUDATarget()) {
+ switch (CurrentTarget()) {
case CUDAFunctionTarget::Global:
case CUDAFunctionTarget::Device:
return SemaDiagnosticBuilder::K_Immediate;
@@ -845,14 +845,14 @@ SemaBase::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfDeviceCode(SourceLocation Lo
return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
}
-Sema::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfHostCode(SourceLocation Loc,
+Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc,
unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
- switch (CurrentCUDATarget()) {
+ switch (CurrentTarget()) {
case CUDAFunctionTarget::Host:
return SemaDiagnosticBuilder::K_Immediate;
case CUDAFunctionTarget::HostDevice:
@@ -874,7 +874,7 @@ Sema::SemaDiagnosticBuilder SemaCUDA::CUDADiagIfHostCode(SourceLocation Loc,
return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
}
-bool SemaCUDA::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
+bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
assert(Callee && "Callee may not be null.");
@@ -894,7 +894,7 @@ bool SemaCUDA::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
SemaRef.getEmissionStatus(Caller) == Sema::FunctionEmissionStatus::Emitted;
SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
CallerKnownEmitted] {
- switch (IdentifyCUDAPreference(Caller, Callee)) {
+ switch (IdentifyPreference(Caller, Callee)) {
case CFP_Never:
case CFP_WrongSide:
assert(Caller && "Never/wrongSide calls require a non-null caller");
@@ -928,8 +928,8 @@ bool SemaCUDA::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
return true;
SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, SemaRef)
- << llvm::to_underlying(IdentifyCUDATarget(Callee)) << /*function*/ 0
- << Callee << llvm::to_underlying(IdentifyCUDATarget(Caller));
+ << llvm::to_underlying(IdentifyTarget(Callee)) << /*function*/ 0
+ << Callee << llvm::to_underlying(IdentifyTarget(Caller));
if (!Callee->getBuiltinID())
SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
diag::note_previous_decl, Caller, SemaRef)
@@ -943,7 +943,7 @@ bool SemaCUDA::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
// defined and uses the capture by reference when the lambda is called. When
// the capture and use happen on different sides, the capture is invalid and
// should be diagnosed.
-void SemaCUDA::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
+void SemaCUDA::CheckLambdaCapture(CXXMethodDecl *Callee,
const sema::Capture &Capture) {
// In host compilation we only need to check lambda functions emitted on host
// side. In such lambda functions, a reference capture is invalid only
@@ -988,7 +988,7 @@ void SemaCUDA::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
}
}
-void SemaCUDA::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
+void SemaCUDA::SetLambdaAttrs(CXXMethodDecl *Method) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
return;
@@ -996,16 +996,16 @@ void SemaCUDA::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
Method->addAttr(CUDAHostAttr::CreateImplicit(getASTContext()));
}
-void SemaCUDA::checkCUDATargetOverload(FunctionDecl *NewFD,
+void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- CUDAFunctionTarget NewTarget = IdentifyCUDATarget(NewFD);
+ CUDAFunctionTarget NewTarget = IdentifyTarget(NewFD);
for (NamedDecl *OldND : Previous) {
FunctionDecl *OldFD = OldND->getAsFunction();
if (!OldFD)
continue;
- CUDAFunctionTarget OldTarget = IdentifyCUDATarget(OldFD);
+ CUDAFunctionTarget OldTarget = IdentifyTarget(OldFD);
// Don't allow HD and global functions to overload other functions with the
// same signature. We allow overloading based on CUDA attributes so that
// functions can have different implementations on the host and device, but
@@ -1014,11 +1014,11 @@ void SemaCUDA::checkCUDATargetOverload(FunctionDecl *NewFD,
if (NewTarget != OldTarget &&
((NewTarget == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
- isCUDAImplicitHostDeviceFunction(NewFD) &&
+ isImplicitHostDeviceFunction(NewFD) &&
OldTarget == CUDAFunctionTarget::Device)) ||
(OldTarget == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
- isCUDAImplicitHostDeviceFunction(OldFD) &&
+ isImplicitHostDeviceFunction(OldFD) &&
NewTarget == CUDAFunctionTarget::Device)) ||
(NewTarget == CUDAFunctionTarget::Global) ||
(OldTarget == CUDAFunctionTarget::Global)) &&
@@ -1044,7 +1044,7 @@ static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
}
}
-void SemaCUDA::inheritCUDATargetAttrs(FunctionDecl *FD,
+void SemaCUDA::inheritTargetAttrs(FunctionDecl *FD,
const FunctionTemplateDecl &TD) {
const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
copyAttrIfPresent<CUDAGlobalAttr>(SemaRef, FD, TemplateFD);
@@ -1052,7 +1052,7 @@ void SemaCUDA::inheritCUDATargetAttrs(FunctionDecl *FD,
copyAttrIfPresent<CUDADeviceAttr>(SemaRef, FD, TemplateFD);
}
-std::string SemaCUDA::getCudaConfigureFuncName() const {
+std::string SemaCUDA::getConfigureFuncName() const {
if (getLangOpts().HIP)
return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration"
: "hipConfigureCall";
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 7d9efb824a96d4..17032d1370521a 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -10596,12 +10596,12 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
// We do not add HD attributes to specializations here because
// they may have different constexpr-ness compared to their
- // templates and, after maybeAddCUDAHostDeviceAttrs() is applied,
+ // templates and, after maybeAddHostDeviceAttrs() is applied,
// may end up with different effective targets. Instead, a
// specialization inherits its target attributes from its template
// in the CheckFunctionTemplateSpecialization() call below.
if (getLangOpts().CUDA && !isFunctionTemplateSpecialization)
- CUDA().maybeAddCUDAHostDeviceAttrs(NewFD, Previous);
+ CUDA().maybeAddHostDeviceAttrs(NewFD, Previous);
// Handle explict specializations of function templates
// and friend function declarations with an explicit
@@ -10899,12 +10899,12 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
if (getLangOpts().CUDA) {
IdentifierInfo *II = NewFD->getIdentifier();
- if (II && II->isStr(CUDA().getCudaConfigureFuncName()) &&
+ if (II && II->isStr(CUDA().getConfigureFuncName()) &&
!NewFD->isInvalidDecl() &&
NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) {
if (!R->castAs<FunctionType>()->getReturnType()->isScalarType())
Diag(NewFD->getLocation(), diag::err_config_scalar_return)
- << CUDA().getCudaConfigureFuncName();
+ << CUDA().getConfigureFuncName();
Context.setcudaConfigureCallDecl(NewFD);
}
@@ -12399,7 +12399,7 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
}
if (!Redeclaration && LangOpts.CUDA)
- CUDA().checkCUDATargetOverload(NewFD, Previous);
+ CUDA().checkTargetOverload(NewFD, Previous);
}
// Check if the function definition uses any AArch64 SME features without
@@ -14416,7 +14416,7 @@ StmtResult Sema::ActOnCXXForRangeIdentifier(Scope *S, SourceLocation IdentLoc,
void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
if (var->isInvalidDecl()) return;
- CUDA().MaybeAddCUDAConstantAttr(var);
+ CUDA().MaybeAddConstantAttr(var);
if (getLangOpts().OpenCL) {
// OpenCL v2.0 s6.12.5 - Every block variable declaration must have an
@@ -14830,7 +14830,7 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) {
// variables whether they are local or not. CUDA also allows
// constant initializers for __constant__ and __device__ variables.
if (getLangOpts().CUDA)
- CUDA().checkAllowedCUDAInitializer(VD);
+ CUDA().checkAllowedInitializer(VD);
// Grab the dllimport or dllexport attribute off of the VarDecl.
const InheritableAttr *DLLAttr = getDLLAttr(VD);
@@ -20667,7 +20667,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
// when compiling for host, device and global functions are never emitted.
// (Technically, we do emit a host-side stub for global functions, but this
// doesn't count for our purposes here.)
- CUDAFunctionTarget T = CUDA().IdentifyCUDATarget(FD);
+ CUDAFunctionTarget T = CUDA().IdentifyTarget(FD);
if (LangOpts.CUDAIsDevice && T == CUDAFunctionTarget::Host)
return FunctionEmissionStatus::CUDADiscarded;
if (!LangOpts.CUDAIsDevice &&
@@ -20692,5 +20692,5 @@ bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {
// for host, only HD functions actually called from the host get marked as
// known-emitted.
return LangOpts.CUDA && !LangOpts.CUDAIsDevice &&
- CUDA().IdentifyCUDATarget(Callee) == CUDAFunctionTarget::Global;
+ CUDA().IdentifyTarget(Callee) == CUDAFunctionTarget::Global;
}
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 0e7e04bca3cf0a..a214ab3beb6e63 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5100,8 +5100,8 @@ static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
return;
}
if (S.getLangOpts().CUDA && VD->hasLocalStorage() &&
- S.CUDA().CUDADiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared)
- << llvm::to_underlying(S.CUDA().CurrentCUDATarget()))
+ S.CUDA().DiagIfHostCode(AL.getLoc(), diag::err_cuda_host_shared)
+ << llvm::to_underlying(S.CUDA().CurrentTarget()))
return;
D->addAttr(::new (S.Context) CUDASharedAttr(S.Context, AL));
}
@@ -5191,7 +5191,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
// in the Decl node for syntactic reasoning, e.g., pretty-printing.
CallingConv CC;
if (S.CheckCallingConvAttr(AL, CC, /*FD*/ nullptr,
- S.CUDA().IdentifyCUDATarget(dyn_cast<FunctionDecl>(D))))
+ S.CUDA().IdentifyTarget(dyn_cast<FunctionDecl>(D))))
return;
if (!isa<ObjCMethodDecl>(D)) {
@@ -5496,7 +5496,7 @@ bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
if (LangOpts.CUDA) {
auto *Aux = Context.getAuxTargetInfo();
assert(FD || CFT != CUDAFunctionTarget::InvalidTarget);
- auto CudaTarget = FD ? CUDA().IdentifyCUDATarget(FD) : CFT;
+ auto CudaTarget = FD ? CUDA().IdentifyTarget(FD) : CFT;
bool CheckHost = false, CheckDevice = false;
switch (CudaTarget) {
case CUDAFunctionTarget::HostDevice:
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index a1e3f639c7d50c..408080662da6ab 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -9877,14 +9877,14 @@ bool Sema::ShouldDeleteSpecialMember(CXXMethodDecl *MD,
// failed.
// For inherited constructors (non-null ICI), CSM may be passed so that MD
// is treated as certain special member, which may not reflect what special
- // member MD really is. However inferCUDATargetForImplicitSpecialMember
+ // member MD really is. However inferTargetForImplicitSpecialMember
// expects CSM to match MD, therefore recalculate CSM.
assert(ICI || CSM == getSpecialMember(MD));
auto RealCSM = CSM;
if (ICI)
RealCSM = getSpecialMember(MD);
- return CUDA().inferCUDATargetForImplicitSpecialMember(RD, RealCSM, MD,
+ return CUDA().inferTargetForImplicitSpecialMember(RD, RealCSM, MD,
SMI.ConstArg, Diagnose);
}
@@ -14056,7 +14056,7 @@ CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
setupImplicitSpecialMemberType(DefaultCon, Context.VoidTy, std::nullopt);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::DefaultConstructor, DefaultCon,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -14342,7 +14342,7 @@ CXXDestructorDecl *Sema::DeclareImplicitDestructor(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(Destructor, Context.VoidTy, std::nullopt);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::Destructor, Destructor,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -14984,7 +14984,7 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(CopyAssignment, RetType, ArgType);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::CopyAssignment, CopyAssignment,
/* ConstRHS */ Const,
/* Diagnose */ false);
@@ -15336,7 +15336,7 @@ CXXMethodDecl *Sema::DeclareImplicitMoveAssignment(CXXRecordDecl *ClassDecl) {
setupImplicitSpecialMemberType(MoveAssignment, RetType, ArgType);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::MoveAssignment, MoveAssignment,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -15734,7 +15734,7 @@ CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
setupImplicitSpecialMemberType(CopyConstructor, Context.VoidTy, ArgType);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::CopyConstructor, CopyConstructor,
/* ConstRHS */ Const,
/* Diagnose */ false);
@@ -15879,7 +15879,7 @@ CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
setupImplicitSpecialMemberType(MoveConstructor, Context.VoidTy, ArgType);
if (getLangOpts().CUDA)
- CUDA().inferCUDATargetForImplicitSpecialMember(
+ CUDA().inferTargetForImplicitSpecialMember(
ClassDecl, CXXSpecialMemberKind::MoveConstructor, MoveConstructor,
/* ConstRHS */ false,
/* Diagnose */ false);
@@ -16185,7 +16185,7 @@ ExprResult Sema::BuildCXXConstructExpr(
DeclInitType->getBaseElementTypeUnsafe()->getAsCXXRecordDecl()) &&
"given constructor for wrong type");
MarkFunctionReferenced(ConstructLoc, Constructor);
- if (getLangOpts().CUDA && !CUDA().CheckCUDACall(ConstructLoc, Constructor))
+ if (getLangOpts().CUDA && !CUDA().CheckCall(ConstructLoc, Constructor))
return ExprError();
return CheckForImmediateInvocation(
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 63f764843ec0e0..823bf36d88bc95 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -309,7 +309,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
DeduceReturnType(FD, Loc))
return true;
- if (getLangOpts().CUDA && !CUDA().CheckCUDACall(Loc, FD))
+ if (getLangOpts().CUDA && !CUDA().CheckCall(Loc, FD))
return true;
}
@@ -17308,7 +17308,7 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc,
// CUDA device code does not support varargs.
if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
if (const FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext)) {
- CUDAFunctionTarget T = CUDA().IdentifyCUDATarget(F);
+ CUDAFunctionTarget T = CUDA().IdentifyTarget(F);
if (T == CUDAFunctionTarget::Global || T == CUDAFunctionTarget::Device ||
T == CUDAFunctionTarget::HostDevice)
return ExprError(Diag(E->getBeginLoc(), diag::err_va_arg_in_device));
@@ -18962,7 +18962,7 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
checkSpecializationReachability(Loc, Func);
if (getLangOpts().CUDA)
- CUDA().CheckCUDACall(Loc, Func);
+ CUDA().CheckCall(Loc, Func);
// If we need a definition, try to create one.
if (NeedDefinition && !Func->getBody()) {
@@ -19109,7 +19109,7 @@ void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
// side. Therefore keep trying until it is recorded.
if (LangOpts.OffloadImplicitHostDeviceTemplates && LangOpts.CUDAIsDevice &&
!getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(Func))
- CUDA().CUDARecordImplicitHostDeviceFuncUsedByDevice(Func);
+ CUDA().RecordImplicitHostDeviceFuncUsedByDevice(Func);
// If this is the first "real" use, act on that.
if (OdrUse == OdrUseContext::Used && !Func->isUsed(/*CheckUsedAttr=*/false)) {
@@ -19182,8 +19182,8 @@ MarkVarDeclODRUsed(ValueDecl *V, SourceLocation Loc, Sema &SemaRef,
if (SemaRef.LangOpts.CUDA && Var->hasGlobalStorage()) {
auto *FD = dyn_cast_or_null<FunctionDecl>(SemaRef.CurContext);
- auto VarTarget = SemaRef.CUDA().IdentifyCUDATarget(Var);
- auto UserTarget = SemaRef.CUDA().IdentifyCUDATarget(FD);
+ auto VarTarget = SemaRef.CUDA().IdentifyTarget(Var);
+ auto UserTarget = SemaRef.CUDA().IdentifyTarget(FD);
if (VarTarget == SemaCUDA::CVT_Host &&
(UserTarget == CUDAFunctionTarget::Device ||
UserTarget == CUDAFunctionTarget::HostDevice ||
diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index 1c80c7467ee1da..1b93e48b4ef4a7 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -885,8 +885,8 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex,
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
- CUDA().CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
- << "throw" << llvm::to_underlying(CUDA().CurrentCUDATarget());
+ CUDA().DiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
+ << "throw" << llvm::to_underlying(CUDA().CurrentTarget());
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw";
@@ -1709,7 +1709,7 @@ bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
// [CUDA] Ignore this function, if we can't call it.
const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
if (getLangOpts().CUDA) {
- auto CallPreference = CUDA().IdentifyCUDAPreference(Caller, Method);
+ auto CallPreference = CUDA().IdentifyPreference(Caller, Method);
// If it's not callable at all, it's not the right function.
if (CallPreference < SemaCUDA::CFP_WrongSide)
return false;
@@ -1719,7 +1719,7 @@ bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
Method->getDeclContext()->lookup(Method->getDeclName());
for (const auto *D : R) {
if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
- if (CUDA().IdentifyCUDAPreference(Caller, FD) > SemaCUDA::CFP_WrongSide)
+ if (CUDA().IdentifyPreference(Caller, FD) > SemaCUDA::CFP_WrongSide)
return false;
}
}
@@ -1738,7 +1738,7 @@ bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) {
return llvm::none_of(PreventedBy, [&](const FunctionDecl *FD) {
assert(FD->getNumParams() == 1 &&
"Only single-operand functions should be in PreventedBy");
- return CUDA().IdentifyCUDAPreference(Caller, FD) >= SemaCUDA::CFP_HostDevice;
+ return CUDA().IdentifyPreference(Caller, FD) >= SemaCUDA::CFP_HostDevice;
});
}
@@ -1801,7 +1801,7 @@ namespace {
// In CUDA, determine how much we'd like / dislike to call this.
if (S.getLangOpts().CUDA)
- CUDAPref = S.CUDA().IdentifyCUDAPreference(
+ CUDAPref = S.CUDA().IdentifyPreference(
S.getCurFunctionDecl(/*AllowLambda=*/true), FD);
}
@@ -2957,7 +2957,7 @@ bool Sema::FindAllocationFunctions(SourceLocation StartLoc, SourceRange Range,
}
if (getLangOpts().CUDA)
- CUDA().EraseUnwantedCUDAMatches(getCurFunctionDecl(/*AllowLambda=*/true),
+ CUDA().EraseUnwantedMatches(getCurFunctionDecl(/*AllowLambda=*/true),
Matches);
} else {
// C++1y [expr.new]p22:
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index 32620675711997..3c38a7e35b093b 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -1394,7 +1394,7 @@ void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
// CUDA lambdas get implicit host and device attributes.
if (getLangOpts().CUDA)
- CUDA().CUDASetLambdaAttrs(Method);
+ CUDA().SetLambdaAttrs(Method);
// OpenMP lambdas might get assumumption attributes.
if (LangOpts.OpenMP)
@@ -2137,7 +2137,7 @@ ExprResult Sema::BuildLambdaExpr(SourceLocation StartLoc, SourceLocation EndLoc,
CaptureInits.push_back(Init.get());
if (LangOpts.CUDA)
- CUDA().CUDACheckLambdaCapture(CallOperator, From);
+ CUDA().CheckLambdaCapture(CallOperator, From);
}
Class->setCaptures(Context, Captures);
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 1b8b3eedfe18fe..6e0847cdcd6079 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -1550,8 +1550,8 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New,
// Don't allow overloading of destructors. (In theory we could, but it
// would be a giant change to clang.)
if (!isa<CXXDestructorDecl>(New)) {
- CUDAFunctionTarget NewTarget = SemaRef.CUDA().IdentifyCUDATarget(New),
- OldTarget = SemaRef.CUDA().IdentifyCUDATarget(Old);
+ CUDAFunctionTarget NewTarget = SemaRef.CUDA().IdentifyTarget(New),
+ OldTarget = SemaRef.CUDA().IdentifyTarget(Old);
if (NewTarget != CUDAFunctionTarget::InvalidTarget) {
assert((OldTarget != CUDAFunctionTarget::InvalidTarget) &&
"Unexpected invalid target.");
@@ -7101,7 +7101,7 @@ void Sema::AddOverloadCandidate(
// inferred for the member automatically, based on the bases and fields of
// the class.
if (!(Caller && Caller->isImplicit()) &&
- !CUDA().IsAllowedCUDACall(Caller, Function)) {
+ !CUDA().IsAllowedCall(Caller, Function)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -7619,7 +7619,7 @@ Sema::AddMethodCandidate(CXXMethodDecl *Method, DeclAccessPair FoundDecl,
// (CUDA B.1): Check for invalid calls between targets.
if (getLangOpts().CUDA)
- if (!CUDA().IsAllowedCUDACall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
+ if (!CUDA().IsAllowedCall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -10441,7 +10441,7 @@ bool clang::isBetterOverloadCandidate(
// If other rules cannot determine which is better, CUDA preference will be
// used again to determine which is better.
//
- // TODO: Currently IdentifyCUDAPreference does not return correct values
+ // TODO: Currently IdentifyPreference does not return correct values
// for functions called in global variable initializers due to missing
// correct context about device/host. Therefore we can only enforce this
// rule when there is a caller. We should enforce this rule for functions
@@ -10453,13 +10453,13 @@ bool clang::isBetterOverloadCandidate(
if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function &&
S.getLangOpts().GPUExcludeWrongSideOverloads) {
if (FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true)) {
- bool IsCallerImplicitHD = SemaCUDA::isCUDAImplicitHostDeviceFunction(Caller);
+ bool IsCallerImplicitHD = SemaCUDA::isImplicitHostDeviceFunction(Caller);
bool IsCand1ImplicitHD =
- SemaCUDA::isCUDAImplicitHostDeviceFunction(Cand1.Function);
+ SemaCUDA::isImplicitHostDeviceFunction(Cand1.Function);
bool IsCand2ImplicitHD =
- SemaCUDA::isCUDAImplicitHostDeviceFunction(Cand2.Function);
- auto P1 = S.CUDA().IdentifyCUDAPreference(Caller, Cand1.Function);
- auto P2 = S.CUDA().IdentifyCUDAPreference(Caller, Cand2.Function);
+ SemaCUDA::isImplicitHostDeviceFunction(Cand2.Function);
+ auto P1 = S.CUDA().IdentifyPreference(Caller, Cand1.Function);
+ auto P2 = S.CUDA().IdentifyPreference(Caller, Cand2.Function);
assert(P1 != SemaCUDA::CFP_Never && P2 != SemaCUDA::CFP_Never);
// The implicit HD function may be a function in a system header which
// is forced by pragma. In device compilation, if we prefer HD candidates
@@ -10759,8 +10759,8 @@ bool clang::isBetterOverloadCandidate(
// to determine which is better.
if (S.getLangOpts().CUDA && Cand1.Function && Cand2.Function) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
- return S.CUDA().IdentifyCUDAPreference(Caller, Cand1.Function) >
- S.CUDA().IdentifyCUDAPreference(Caller, Cand2.Function);
+ return S.CUDA().IdentifyPreference(Caller, Cand1.Function) >
+ S.CUDA().IdentifyPreference(Caller, Cand2.Function);
}
// General member function overloading is handled above, so this only handles
@@ -10892,14 +10892,14 @@ OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
// Check viable function only.
return Cand->Viable && Cand->Function &&
- S.CUDA().IdentifyCUDAPreference(Caller, Cand->Function) ==
+ S.CUDA().IdentifyPreference(Caller, Cand->Function) ==
SemaCUDA::CFP_SameSide;
});
if (ContainsSameSideCandidate) {
auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {
// Check viable function only to avoid unnecessary data copying/moving.
return Cand->Viable && Cand->Function &&
- S.CUDA().IdentifyCUDAPreference(Caller, Cand->Function) ==
+ S.CUDA().IdentifyPreference(Caller, Cand->Function) ==
SemaCUDA::CFP_WrongSide;
};
llvm::erase_if(Candidates, IsWrongSideCandidate);
@@ -11939,8 +11939,8 @@ static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
FunctionDecl *Callee = Cand->Function;
- CUDAFunctionTarget CallerTarget = S.CUDA().IdentifyCUDATarget(Caller),
- CalleeTarget = S.CUDA().IdentifyCUDATarget(Callee);
+ CUDAFunctionTarget CallerTarget = S.CUDA().IdentifyTarget(Caller),
+ CalleeTarget = S.CUDA().IdentifyTarget(Callee);
std::string FnDesc;
std::pair<OverloadCandidateKind, OverloadCandidateSelect> FnKindPair =
@@ -11987,7 +11987,7 @@ static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
}
}
- S.CUDA().inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth,
+ S.CUDA().inferTargetForImplicitSpecialMember(ParentClass, CSM, Meth,
/* ConstRHS */ ConstRHS,
/* Diagnose */ true);
}
@@ -13061,7 +13061,7 @@ class AddressOfFunctionResolver {
if (S.getLangOpts().CUDA) {
FunctionDecl *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true);
if (!(Caller && Caller->isImplicit()) &&
- !S.CUDA().IsAllowedCUDACall(Caller, FunDecl))
+ !S.CUDA().IsAllowedCall(Caller, FunDecl))
return false;
}
if (FunDecl->isMultiVersion()) {
@@ -13181,7 +13181,7 @@ class AddressOfFunctionResolver {
}
void EliminateSuboptimalCudaMatches() {
- S.CUDA().EraseUnwantedCUDAMatches(S.getCurFunctionDecl(/*AllowLambda=*/true),
+ S.CUDA().EraseUnwantedMatches(S.getCurFunctionDecl(/*AllowLambda=*/true),
Matches);
}
@@ -13336,8 +13336,8 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
// Return positive for better, negative for worse, 0 for equal preference.
auto CheckCUDAPreference = [&](FunctionDecl *FD1, FunctionDecl *FD2) {
FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
- return static_cast<int>(CUDA().IdentifyCUDAPreference(Caller, FD1)) -
- static_cast<int>(CUDA().IdentifyCUDAPreference(Caller, FD2));
+ return static_cast<int>(CUDA().IdentifyPreference(Caller, FD1)) -
+ static_cast<int>(CUDA().IdentifyPreference(Caller, FD2));
};
auto CheckMoreConstrained = [&](FunctionDecl *FD1,
diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 732e34257c47d8..d28c24cfdfd33c 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -4575,8 +4575,8 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
// Exceptions aren't allowed in CUDA device code.
if (getLangOpts().CUDA)
- CUDA().CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
- << "try" << llvm::to_underlying(CUDA().CurrentCUDATarget());
+ CUDA().DiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
+ << "try" << llvm::to_underlying(CUDA().CurrentTarget());
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try";
diff --git a/clang/lib/Sema/SemaTemplate.cpp b/clang/lib/Sema/SemaTemplate.cpp
index fa1127b10aa7e2..3b9f8e5d638714 100644
--- a/clang/lib/Sema/SemaTemplate.cpp
+++ b/clang/lib/Sema/SemaTemplate.cpp
@@ -10156,9 +10156,9 @@ bool Sema::CheckFunctionTemplateSpecialization(
// take target attributes into account, we reject candidates
// here that have a different target.
if (LangOpts.CUDA &&
- CUDA().IdentifyCUDATarget(Specialization,
+ CUDA().IdentifyTarget(Specialization,
/* IgnoreImplicitHDAttr = */ true) !=
- CUDA().IdentifyCUDATarget(FD, /* IgnoreImplicitHDAttr = */ true)) {
+ CUDA().IdentifyTarget(FD, /* IgnoreImplicitHDAttr = */ true)) {
FailedCandidates.addCandidate().set(
I.getPair(), FunTmpl->getTemplatedDecl(),
MakeDeductionFailureInfo(
@@ -10329,7 +10329,7 @@ bool Sema::CheckFunctionTemplateSpecialization(
// virtue e.g. of being constexpr, and it passes these implicit
// attributes on to its specializations.)
if (LangOpts.CUDA)
- CUDA().inheritCUDATargetAttrs(FD, *Specialization->getPrimaryTemplate());
+ CUDA().inheritTargetAttrs(FD, *Specialization->getPrimaryTemplate());
// The "previous declaration" for this function template specialization is
// the prior function template specialization.
@@ -11365,9 +11365,9 @@ DeclResult Sema::ActOnExplicitInstantiation(Scope *S,
// target attributes into account, we reject candidates here that
// have a different target.
if (LangOpts.CUDA &&
- CUDA().IdentifyCUDATarget(Specialization,
+ CUDA().IdentifyTarget(Specialization,
/* IgnoreImplicitHDAttr = */ true) !=
- CUDA().IdentifyCUDATarget(D.getDeclSpec().getAttributes())) {
+ CUDA().IdentifyTarget(D.getDeclSpec().getAttributes())) {
FailedCandidates.addCandidate().set(
P.getPair(), FunTmpl->getTemplatedDecl(),
MakeDeductionFailureInfo(
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 15e247a956e3ed..c0469a47ab8b62 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -5538,7 +5538,7 @@ void Sema::InstantiateVariableInitializer(
}
if (getLangOpts().CUDA)
- CUDA().checkAllowedCUDAInitializer(Var);
+ CUDA().checkAllowedInitializer(Var);
}
/// Instantiate the definition of the given variable from its
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 09769ced5d9c21..404c4e8e31b558 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2736,7 +2736,7 @@ QualType Sema::BuildArrayType(QualType T, ArraySizeModifier ASM,
bool IsCUDADevice = (getLangOpts().CUDA && getLangOpts().CUDAIsDevice);
targetDiag(Loc,
IsCUDADevice ? diag::err_cuda_vla : diag::err_vla_unsupported)
- << (IsCUDADevice ? llvm::to_underlying(CUDA().CurrentCUDATarget()) : 0);
+ << (IsCUDADevice ? llvm::to_underlying(CUDA().CurrentTarget()) : 0);
} else if (sema::FunctionScopeInfo *FSI = getCurFunction()) {
// VLAs are supported on this target, but we may need to do delayed
// checking that the VLA is not being used within a coroutine.
@@ -3619,7 +3619,7 @@ static QualType GetDeclSpecTypeForDeclarator(TypeProcessingState &state,
// D.getDeclarationAttributes()) because those are always C++11 attributes,
// and those don't get distributed.
distributeTypeAttrsFromDeclarator(
- state, T, SemaRef.CUDA().IdentifyCUDATarget(D.getAttributes()));
+ state, T, SemaRef.CUDA().IdentifyTarget(D.getAttributes()));
// Find the deduced type in this type. Look in the trailing return type if we
// have one, otherwise in the DeclSpec type.
@@ -4140,7 +4140,7 @@ static CallingConv getCCForDeclaratorChunk(
// handleFunctionTypeAttr.
CallingConv CC;
if (!S.CheckCallingConvAttr(AL, CC, /*FunctionDecl=*/nullptr,
- S.CUDA().IdentifyCUDATarget(D.getAttributes())) &&
+ S.CUDA().IdentifyTarget(D.getAttributes())) &&
(!FTI.isVariadic || supportsVariadicCall(CC))) {
return CC;
}
@@ -5826,7 +5826,7 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
// See if there are any attributes on this declarator chunk.
processTypeAttrs(state, T, TAL_DeclChunk, DeclType.getAttrs(),
- S.CUDA().IdentifyCUDATarget(D.getAttributes()));
+ S.CUDA().IdentifyTarget(D.getAttributes()));
if (DeclType.Kind != DeclaratorChunk::Paren) {
if (ExpectNoDerefChunk && !IsNoDerefableChunk(DeclType))
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 93051d05cc4e8c..c4688847483870 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -3996,7 +3996,7 @@ llvm::Error ASTReader::ReadASTBlock(ModuleFile &F,
if (Record.size() != 1)
return llvm::createStringError(std::errc::illegal_byte_sequence,
"invalid cuda pragma options record");
- ForceCUDAHostDeviceDepth = Record[0];
+ ForceHostDeviceDepth = Record[0];
break;
case ALIGN_PACK_PRAGMA_OPTIONS: {
@@ -8275,7 +8275,7 @@ void ASTReader::UpdateSema() {
PragmaMSPointersToMembersState,
PointersToMembersPragmaLocation);
}
- SemaObj->CUDA().ForceCUDAHostDeviceDepth = ForceCUDAHostDeviceDepth;
+ SemaObj->CUDA().ForceHostDeviceDepth = ForceHostDeviceDepth;
if (PragmaAlignPackCurrentValue) {
// The bottom of the stack might have a default value. It must be adjusted
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index d959bb5e69d03b..88f93feaf5cb05 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -4336,8 +4336,8 @@ void ASTWriter::WriteOpenCLExtensions(Sema &SemaRef) {
Stream.EmitRecord(OPENCL_EXTENSIONS, Record);
}
void ASTWriter::WriteCUDAPragmas(Sema &SemaRef) {
- if (SemaRef.CUDA().ForceCUDAHostDeviceDepth > 0) {
- RecordData::value_type Record[] = {SemaRef.CUDA().ForceCUDAHostDeviceDepth};
+ if (SemaRef.CUDA().ForceHostDeviceDepth > 0) {
+ RecordData::value_type Record[] = {SemaRef.CUDA().ForceHostDeviceDepth};
Stream.EmitRecord(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH, Record);
}
}
>From b33cc26fde62b89625cee590ee9c9fabdb949200 Mon Sep 17 00:00:00 2001
From: Vlad Serebrennikov <serebrennikov.vladislav at gmail.com>
Date: Fri, 12 Apr 2024 21:47:36 +0300
Subject: [PATCH 3/4] Run clang-format
---
clang/include/clang/Sema/SemaCUDA.h | 32 ++++++-------
clang/lib/Parse/ParseDecl.cpp | 3 +-
clang/lib/Parse/ParseExpr.cpp | 6 +--
clang/lib/Sema/SemaCUDA.cpp | 72 ++++++++++++++++-------------
clang/lib/Sema/SemaDeclAttr.cpp | 5 +-
clang/lib/Sema/SemaDeclCXX.cpp | 2 +-
clang/lib/Sema/SemaExprCXX.cpp | 2 +-
clang/lib/Sema/SemaLambda.cpp | 4 +-
clang/lib/Sema/SemaOverload.cpp | 9 ++--
clang/lib/Sema/SemaTemplate.cpp | 4 +-
10 files changed, 72 insertions(+), 67 deletions(-)
diff --git a/clang/include/clang/Sema/SemaCUDA.h b/clang/include/clang/Sema/SemaCUDA.h
index b7dd57aaeadcd0..71cde5a49f6b1a 100644
--- a/clang/include/clang/Sema/SemaCUDA.h
+++ b/clang/include/clang/Sema/SemaCUDA.h
@@ -47,8 +47,8 @@ class SemaCUDA : public SemaBase {
bool PopForceHostDevice();
ExprResult ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
- MultiExprArg ExecConfig,
- SourceLocation GGGLoc);
+ 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.
@@ -88,8 +88,7 @@ class SemaCUDA : public SemaBase {
/// if (DiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentTarget())
/// return ExprError();
/// // Otherwise, continue parsing as normal.
- SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc,
- unsigned DiagID);
+ SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
/// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
/// context is "used as host code".
@@ -103,7 +102,7 @@ class SemaCUDA : public SemaBase {
/// 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);
+ bool IgnoreImplicitHDAttr = false);
CUDAFunctionTarget IdentifyTarget(const ParsedAttributesView &Attrs);
enum CUDAVariableTarget {
@@ -134,7 +133,8 @@ class SemaCUDA : public SemaBase {
struct CUDATargetContextRAII {
SemaCUDA &S;
SemaCUDA::CUDATargetContext SavedCtx;
- CUDATargetContextRAII(SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K, Decl *D);
+ CUDATargetContextRAII(SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K,
+ Decl *D);
~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
};
@@ -166,22 +166,20 @@ class SemaCUDA : public SemaBase {
///
/// \returns preference value for particular Caller/Callee combination.
CUDAFunctionPreference IdentifyPreference(const FunctionDecl *Caller,
- const FunctionDecl *Callee);
+ 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) {
+ 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);
+ void maybeAddHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous);
/// May add implicit CUDAConstantAttr attribute to VD, depending on VD
/// and current compilation settings.
@@ -221,7 +219,8 @@ class SemaCUDA : public SemaBase {
/// calling priority.
void EraseUnwantedMatches(
const FunctionDecl *Caller,
- llvm::SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches);
+ 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.
@@ -235,9 +234,9 @@ class SemaCUDA : public SemaBase {
/// 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);
+ 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).
@@ -256,8 +255,7 @@ class SemaCUDA : public SemaBase {
/// Check whether NewFD is a valid overload for CUDA. Emits
/// diagnostics and invalidates NewFD if not.
- void checkTargetOverload(FunctionDecl *NewFD,
- const LookupResult &Previous);
+ 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);
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 37be68d6ec5be2..951e2210031a14 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -2665,7 +2665,8 @@ Decl *Parser::ParseDeclarationAfterDeclaratorAndAttributes(
}
}
- SemaCUDA::CUDATargetContextRAII X(Actions.CUDA(), SemaCUDA::CTCK_InitGlobalVar, ThisDecl);
+ SemaCUDA::CUDATargetContextRAII X(Actions.CUDA(),
+ SemaCUDA::CTCK_InitGlobalVar, ThisDecl);
switch (TheInitKind) {
// Parse declarator '=' initializer.
case InitKind::Equal: {
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index ec05ce616ca647..473ec9afd60181 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -2130,10 +2130,8 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
}
if (!LHS.isInvalid()) {
- ExprResult ECResult = Actions.CUDA().ActOnExecConfigExpr(getCurScope(),
- OpenLoc,
- ExecConfigExprs,
- CloseLoc);
+ ExprResult ECResult = Actions.CUDA().ActOnExecConfigExpr(
+ getCurScope(), OpenLoc, ExecConfigExprs, CloseLoc);
if (ECResult.isInvalid())
LHS = ExprError();
else
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13d318fe911766..80ea43dc5316eb 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -60,12 +60,12 @@ ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
<< getConfigureFuncName());
QualType ConfigQTy = ConfigDecl->getType();
- DeclRefExpr *ConfigDR = new (getASTContext())
- DeclRefExpr(getASTContext(), ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
+ DeclRefExpr *ConfigDR = new (getASTContext()) DeclRefExpr(
+ getASTContext(), ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
SemaRef.MarkFunctionReferenced(LLLLoc, ConfigDecl);
return SemaRef.BuildCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
- /*IsExecConfig=*/true);
+ /*IsExecConfig=*/true);
}
CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
@@ -115,9 +115,8 @@ static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
});
}
-SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(SemaCUDA &S_,
- SemaCUDA::CUDATargetContextKind K,
- Decl *D)
+SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(
+ SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K, Decl *D)
: S(S_) {
SavedCtx = S.CurCUDATargetCtx;
assert(K == SemaCUDA::CTCK_InitGlobalVar);
@@ -416,11 +415,11 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
Sema::SpecialMemberOverloadResult SMOR =
SemaRef.LookupSpecialMember(BaseClassDecl, CSM,
- /* ConstArg */ ConstRHS,
- /* VolatileArg */ false,
- /* RValueThis */ false,
- /* ConstThis */ false,
- /* VolatileThis */ false);
+ /* ConstArg */ ConstRHS,
+ /* VolatileArg */ false,
+ /* RValueThis */ false,
+ /* ConstThis */ false,
+ /* VolatileThis */ false);
if (!SMOR.getMethod())
continue;
@@ -438,7 +437,8 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
<< (unsigned)CSM << llvm::to_underlying(*InferredTarget)
<< llvm::to_underlying(BaseMethodTarget);
}
- MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
+ MemberDecl->addAttr(
+ CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
return true;
}
}
@@ -459,17 +459,16 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
Sema::SpecialMemberOverloadResult SMOR =
SemaRef.LookupSpecialMember(FieldRecDecl, CSM,
- /* ConstArg */ ConstRHS && !F->isMutable(),
- /* VolatileArg */ false,
- /* RValueThis */ false,
- /* ConstThis */ false,
- /* VolatileThis */ false);
+ /* ConstArg */ ConstRHS && !F->isMutable(),
+ /* VolatileArg */ false,
+ /* RValueThis */ false,
+ /* ConstThis */ false,
+ /* VolatileThis */ false);
if (!SMOR.getMethod())
continue;
- CUDAFunctionTarget FieldMethodTarget =
- IdentifyTarget(SMOR.getMethod());
+ CUDAFunctionTarget FieldMethodTarget = IdentifyTarget(SMOR.getMethod());
if (!InferredTarget) {
InferredTarget = FieldMethodTarget;
} else {
@@ -482,7 +481,8 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
<< (unsigned)CSM << llvm::to_underlying(*InferredTarget)
<< llvm::to_underlying(FieldMethodTarget);
}
- MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
+ MemberDecl->addAttr(
+ CUDAInvalidTargetAttr::CreateImplicit(getASTContext()));
return true;
}
}
@@ -774,8 +774,9 @@ void SemaCUDA::maybeAddHostDeviceAttrs(FunctionDecl *NewD,
FunctionDecl *OldD = D->getAsFunction();
return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
!OldD->hasAttr<CUDAHostAttr>() &&
- !SemaRef.IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
- /* ConsiderCudaAttrs = */ false);
+ !SemaRef.IsOverload(NewD, OldD,
+ /* UseMemberUsingDeclRules = */ false,
+ /* ConsiderCudaAttrs = */ false);
};
auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
if (It != Previous.end()) {
@@ -816,9 +817,10 @@ void SemaCUDA::MaybeAddConstantAttr(VarDecl *VD) {
}
SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
- unsigned DiagID) {
+ unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *CurFunContext =
+ SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
@@ -832,7 +834,8 @@ SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
// mode until the function is known-emitted.
if (!getLangOpts().CUDAIsDevice)
return SemaDiagnosticBuilder::K_Nop;
- if (SemaRef.IsLastErrorImmediate && getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
+ if (SemaRef.IsLastErrorImmediate &&
+ getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
return SemaDiagnosticBuilder::K_Immediate;
return (SemaRef.getEmissionStatus(CurFunContext) ==
Sema::FunctionEmissionStatus::Emitted)
@@ -848,7 +851,8 @@ SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc,
unsigned DiagID) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
- FunctionDecl *CurFunContext = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
+ FunctionDecl *CurFunContext =
+ SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
SemaDiagnosticBuilder::Kind DiagKind = [&] {
if (!CurFunContext)
return SemaDiagnosticBuilder::K_Nop;
@@ -861,7 +865,8 @@ Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc,
// mode until the function is known-emitted.
if (getLangOpts().CUDAIsDevice)
return SemaDiagnosticBuilder::K_Nop;
- if (SemaRef.IsLastErrorImmediate && getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
+ if (SemaRef.IsLastErrorImmediate &&
+ getDiagnostics().getDiagnosticIDs()->isBuiltinNote(DiagID))
return SemaDiagnosticBuilder::K_Immediate;
return (SemaRef.getEmissionStatus(CurFunContext) ==
Sema::FunctionEmissionStatus::Emitted)
@@ -890,8 +895,8 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
// If the caller is known-emitted, mark the callee as known-emitted.
// Otherwise, mark the call in our call graph so we can traverse it later.
- bool CallerKnownEmitted =
- SemaRef.getEmissionStatus(Caller) == Sema::FunctionEmissionStatus::Emitted;
+ bool CallerKnownEmitted = SemaRef.getEmissionStatus(Caller) ==
+ Sema::FunctionEmissionStatus::Emitted;
SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
CallerKnownEmitted] {
switch (IdentifyPreference(Caller, Callee)) {
@@ -927,9 +932,10 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
return true;
- SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, SemaRef)
- << llvm::to_underlying(IdentifyTarget(Callee)) << /*function*/ 0
- << Callee << llvm::to_underlying(IdentifyTarget(Caller));
+ SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller,
+ SemaRef)
+ << llvm::to_underlying(IdentifyTarget(Callee)) << /*function*/ 0 << Callee
+ << llvm::to_underlying(IdentifyTarget(Caller));
if (!Callee->getBuiltinID())
SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
diag::note_previous_decl, Caller, SemaRef)
@@ -1023,7 +1029,7 @@ void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
(NewTarget == CUDAFunctionTarget::Global) ||
(OldTarget == CUDAFunctionTarget::Global)) &&
!SemaRef.IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
- /* ConsiderCudaAttrs = */ false)) {
+ /* ConsiderCudaAttrs = */ false)) {
Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
<< llvm::to_underlying(NewTarget) << NewFD->getDeclName()
<< llvm::to_underlying(OldTarget) << OldFD;
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index a214ab3beb6e63..b7b1fbc625a150 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5190,8 +5190,9 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
// Diagnostic is emitted elsewhere: here we store the (valid) AL
// in the Decl node for syntactic reasoning, e.g., pretty-printing.
CallingConv CC;
- if (S.CheckCallingConvAttr(AL, CC, /*FD*/ nullptr,
- S.CUDA().IdentifyTarget(dyn_cast<FunctionDecl>(D))))
+ if (S.CheckCallingConvAttr(
+ AL, CC, /*FD*/ nullptr,
+ S.CUDA().IdentifyTarget(dyn_cast<FunctionDecl>(D))))
return;
if (!isa<ObjCMethodDecl>(D)) {
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index 408080662da6ab..1fe10375222c53 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -9885,7 +9885,7 @@ bool Sema::ShouldDeleteSpecialMember(CXXMethodDecl *MD,
RealCSM = getSpecialMember(MD);
return CUDA().inferTargetForImplicitSpecialMember(RD, RealCSM, MD,
- SMI.ConstArg, Diagnose);
+ SMI.ConstArg, Diagnose);
}
return false;
diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index 1b93e48b4ef4a7..8911257a6f6146 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -2958,7 +2958,7 @@ bool Sema::FindAllocationFunctions(SourceLocation StartLoc, SourceRange Range,
if (getLangOpts().CUDA)
CUDA().EraseUnwantedMatches(getCurFunctionDecl(/*AllowLambda=*/true),
- Matches);
+ Matches);
} else {
// C++1y [expr.new]p22:
// For a non-placement allocation function, the normal deallocation
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index 3c38a7e35b093b..35a51c6c2328db 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -9,18 +9,18 @@
// This file implements semantic analysis for C++ lambda expressions.
//
//===----------------------------------------------------------------------===//
-#include "clang/Sema/DeclSpec.h"
+#include "clang/Sema/SemaLambda.h"
#include "TypeLocBuilder.h"
#include "clang/AST/ASTLambda.h"
#include "clang/AST/ExprCXX.h"
#include "clang/Basic/TargetInfo.h"
+#include "clang/Sema/DeclSpec.h"
#include "clang/Sema/Initialization.h"
#include "clang/Sema/Lookup.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ScopeInfo.h"
#include "clang/Sema/SemaCUDA.h"
#include "clang/Sema/SemaInternal.h"
-#include "clang/Sema/SemaLambda.h"
#include "clang/Sema/Template.h"
#include "llvm/ADT/STLExtras.h"
#include <optional>
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 6e0847cdcd6079..397e7681828f39 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -7619,7 +7619,8 @@ Sema::AddMethodCandidate(CXXMethodDecl *Method, DeclAccessPair FoundDecl,
// (CUDA B.1): Check for invalid calls between targets.
if (getLangOpts().CUDA)
- if (!CUDA().IsAllowedCall(getCurFunctionDecl(/*AllowLambda=*/true), Method)) {
+ if (!CUDA().IsAllowedCall(getCurFunctionDecl(/*AllowLambda=*/true),
+ Method)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -11988,8 +11989,8 @@ static void DiagnoseBadTarget(Sema &S, OverloadCandidate *Cand) {
}
S.CUDA().inferTargetForImplicitSpecialMember(ParentClass, CSM, Meth,
- /* ConstRHS */ ConstRHS,
- /* Diagnose */ true);
+ /* ConstRHS */ ConstRHS,
+ /* Diagnose */ true);
}
}
@@ -13182,7 +13183,7 @@ class AddressOfFunctionResolver {
void EliminateSuboptimalCudaMatches() {
S.CUDA().EraseUnwantedMatches(S.getCurFunctionDecl(/*AllowLambda=*/true),
- Matches);
+ Matches);
}
public:
diff --git a/clang/lib/Sema/SemaTemplate.cpp b/clang/lib/Sema/SemaTemplate.cpp
index 3b9f8e5d638714..95171359f0ab17 100644
--- a/clang/lib/Sema/SemaTemplate.cpp
+++ b/clang/lib/Sema/SemaTemplate.cpp
@@ -10157,7 +10157,7 @@ bool Sema::CheckFunctionTemplateSpecialization(
// here that have a different target.
if (LangOpts.CUDA &&
CUDA().IdentifyTarget(Specialization,
- /* IgnoreImplicitHDAttr = */ true) !=
+ /* IgnoreImplicitHDAttr = */ true) !=
CUDA().IdentifyTarget(FD, /* IgnoreImplicitHDAttr = */ true)) {
FailedCandidates.addCandidate().set(
I.getPair(), FunTmpl->getTemplatedDecl(),
@@ -11366,7 +11366,7 @@ DeclResult Sema::ActOnExplicitInstantiation(Scope *S,
// have a different target.
if (LangOpts.CUDA &&
CUDA().IdentifyTarget(Specialization,
- /* IgnoreImplicitHDAttr = */ true) !=
+ /* IgnoreImplicitHDAttr = */ true) !=
CUDA().IdentifyTarget(D.getDeclSpec().getAttributes())) {
FailedCandidates.addCandidate().set(
P.getPair(), FunTmpl->getTemplatedDecl(),
>From 9baa1c9ced65ca391156d2b5866f1a98e607b4df Mon Sep 17 00:00:00 2001
From: Vlad Serebrennikov <serebrennikov.vladislav at gmail.com>
Date: Fri, 12 Apr 2024 21:50:30 +0300
Subject: [PATCH 4/4] Remove an entry from the table of contents in `Sema`
---
clang/include/clang/Sema/Sema.h | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 5dac8083714b61..6b9789334811ec 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -480,8 +480,7 @@ class Sema final : public SemaBase {
// 35. Code Completion (SemaCodeComplete.cpp)
// 36. FixIt Helpers (SemaFixItUtils.cpp)
// 37. Name Lookup for RISC-V Vector Intrinsic (SemaRISCVVectorLookup.cpp)
- // 38. CUDA (SemaCUDA.cpp)
- // 39. OpenMP Directives and Clauses (SemaOpenMP.cpp)
+ // 38. OpenMP Directives and Clauses (SemaOpenMP.cpp)
/// \name Semantic Analysis
/// Implementations are in Sema.cpp
More information about the cfe-commits
mailing list