SemaCUDA.cpp | Explore in Territory

//===--- SemaCUDA.cpp - 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 implements semantic analysis for CUDA constructs.
///
//===----------------------------------------------------------------------===//

#include "clang/Sema/SemaCUDA.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Decl.h"
#include "clang/AST/ExprCXX.h"
#include "clang/Basic/Cuda.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Lex/Preprocessor.h"
#include "clang/Sema/Lookup.h"
#include "clang/Sema/ScopeInfo.h"
#include "clang/Sema/Sema.h"
#include "clang/Sema/SemaDiagnostic.h"
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
#include "llvm/ADT/STLForwardCompat.h"
#include "llvm/ADT/SmallVector.h"
#include <optional>
usingnamespaceclang;

SemaCUDA::SemaCUDA(Sema &S) : … { … }

template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) { … }

void SemaCUDA::PushForceHostDevice() { … }

bool SemaCUDA::PopForceHostDevice() { … }

ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                         MultiExprArg ExecConfig,
                                         SourceLocation GGGLoc) { … }

CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) { … }

template <typename A>
static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) { … }

SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(
    SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K, Decl *D)
    : … { … }

/// IdentifyTarget - Determine the CUDA compilation target for this function
CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
                                            bool IgnoreImplicitHDAttr) { … }

/// IdentifyTarget - Determine the CUDA compilation target for this variable.
SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyTarget(const VarDecl *Var) { … }

// * CUDA Call preference table
//
// F - from,
// T - to
// Ph - preference in host mode
// Pd - preference in device mode
// H  - handled in (x)
// Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.
//
// | F  | T  | Ph  | Pd  |  H  |
// |----+----+-----+-----+-----+
// | d  | d  | N   | N   | (c) |
// | d  | g  | --  | --  | (a) |
// | d  | h  | --  | --  | (e) |
// | d  | hd | HD  | HD  | (b) |
// | g  | d  | N   | N   | (c) |
// | g  | g  | --  | --  | (a) |
// | g  | h  | --  | --  | (e) |
// | g  | hd | HD  | HD  | (b) |
// | h  | d  | --  | --  | (e) |
// | h  | g  | N   | N   | (c) |
// | h  | h  | N   | N   | (c) |
// | h  | hd | HD  | HD  | (b) |
// | hd | d  | WS  | SS  | (d) |
// | hd | g  | SS  | --  |(d/a)|
// | hd | h  | SS  | WS  | (d) |
// | hd | hd | HD  | HD  | (b) |

SemaCUDA::CUDAFunctionPreference
SemaCUDA::IdentifyPreference(const FunctionDecl *Caller,
                             const FunctionDecl *Callee) { … }

template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) { … }

bool SemaCUDA::isImplicitHostDeviceFunction(const FunctionDecl *D) { … }

void SemaCUDA::EraseUnwantedMatches(
    const FunctionDecl *Caller,
    SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) { … }

/// When an implicitly-declared special member has to invoke more than one
/// base/field special member, conflicts may occur in the targets of these
/// members. For example, if one base's member __host__ and another's is
/// __device__, it's a conflict.
/// This function figures out if the given targets \param Target1 and
/// \param Target2 conflict, and if they do not it fills in
/// \param ResolvedTarget with a target that resolves for both calls.
/// \return true if there's a conflict, false otherwise.
static bool
resolveCalleeCUDATargetConflict(CUDAFunctionTarget Target1,
                                CUDAFunctionTarget Target2,
                                CUDAFunctionTarget *ResolvedTarget) { … }

bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                                   CXXSpecialMemberKind CSM,
                                                   CXXMethodDecl *MemberDecl,
                                                   bool ConstRHS,
                                                   bool Diagnose) { … }

bool SemaCUDA::isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD) { … }

bool SemaCUDA::isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *DD) { … }

namespace {
enum CUDAInitializerCheckKind { … };

bool IsDependentVar(VarDecl *VD) { … }

// Check whether a variable has an allowed initializer for a CUDA device side
// variable with global storage. \p VD may be a host variable to be checked for
// potential promotion to device side variable.
//
// CUDA/HIP 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.
bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
                                           CUDAInitializerCheckKind CheckKind) { … }
} // namespace

void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { … }

void SemaCUDA::RecordImplicitHostDeviceFuncUsedByDevice(
    const FunctionDecl *Callee) { … }

// With -fcuda-host-device-constexpr, an unattributed constexpr function is
// treated as implicitly __host__ __device__, unless:
//  * it is a variadic function (device-side variadic functions are not
//    allowed), or
//  * a __device__ function with this signature was already declared, in which
//    case in which case we output an error, unless the __device__ decl is in a
//    system header, in which case we leave the constexpr function unattributed.
//
// In addition, all function decls are treated as __host__ __device__ when
// ForceHostDeviceDepth > 0 (corresponding to code within a
//   #pragma clang force_cuda_host_device_begin/end
// pair).
void SemaCUDA::maybeAddHostDeviceAttrs(FunctionDecl *NewD,
                                       const LookupResult &Previous) { … }

// 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::MaybeAddConstantAttr(VarDecl *VD) { … }

SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
                                                           unsigned DiagID) { … }

Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc,
                                                     unsigned DiagID) { … }

bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) { … }

// Check the wrong-sided reference capture of lambda for CUDA/HIP.
// A lambda function may capture a stack variable by reference when it is
// 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::CheckLambdaCapture(CXXMethodDecl *Callee,
                                  const sema::Capture &Capture) { … }

void SemaCUDA::SetLambdaAttrs(CXXMethodDecl *Method) { … }

void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
                                   const LookupResult &Previous) { … }

template <typename AttrTy>
static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
                              const FunctionDecl &TemplateFD) { … }

void SemaCUDA::inheritTargetAttrs(FunctionDecl *FD,
                                  const FunctionTemplateDecl &TD) { … }

std::string SemaCUDA::getConfigureFuncName() const { … }
llvm/clang/lib/Sema/SemaCUDA.cpp
