clang 20.0.0git
SemaCUDA.h
Go to the documentation of this file.
1//===----- SemaCUDA.h ----- Semantic Analysis for CUDA constructs ---------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8/// \file
9/// This file declares semantic analysis for CUDA constructs.
10///
11//===----------------------------------------------------------------------===//
12
13#ifndef LLVM_CLANG_SEMA_SEMACUDA_H
14#define LLVM_CLANG_SEMA_SEMACUDA_H
15
16#include "clang/AST/Decl.h"
17#include "clang/AST/DeclCXX.h"
19#include "clang/Basic/Cuda.h"
21#include "clang/Sema/Lookup.h"
24#include "clang/Sema/Scope.h"
26#include "clang/Sema/SemaBase.h"
27#include "llvm/ADT/DenseMap.h"
28#include "llvm/ADT/SmallVector.h"
29#include <string>
30
31namespace clang {
32
33enum class CUDAFunctionTarget;
34
35class SemaCUDA : public SemaBase {
36public:
37 SemaCUDA(Sema &S);
38
39 /// Increments our count of the number of times we've seen a pragma forcing
40 /// functions to be __host__ __device__. So long as this count is greater
41 /// than zero, all functions encountered will be __host__ __device__.
43
44 /// Decrements our count of the number of times we've seen a pragma forcing
45 /// functions to be __host__ __device__. Returns false if the count is 0
46 /// before incrementing, so you can emit an error.
47 bool PopForceHostDevice();
48
50 MultiExprArg ExecConfig,
51 SourceLocation GGGLoc);
52
53 /// A pair of a canonical FunctionDecl and a SourceLocation. When used as the
54 /// key in a hashtable, both the FD and location are hashed.
58 };
59
60 /// FunctionDecls and SourceLocations for which CheckCall has emitted a
61 /// (maybe deferred) "bad call" diagnostic. We use this to avoid emitting the
62 /// same deferred diag twice.
64
65 /// An inverse call graph, mapping known-emitted functions to one of their
66 /// known-emitted callers (plus the location of the call).
67 ///
68 /// Functions that we can tell a priori must be emitted aren't added to this
69 /// map.
70 llvm::DenseMap</* Callee = */ CanonicalDeclPtr<const FunctionDecl>,
71 /* Caller = */ FunctionDeclAndLoc>
73
74 /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
75 /// context is "used as device code".
76 ///
77 /// - If CurContext is a __host__ function, does not emit any diagnostics
78 /// unless \p EmitOnBothSides is true.
79 /// - If CurContext is a __device__ or __global__ function, emits the
80 /// diagnostics immediately.
81 /// - If CurContext is a __host__ __device__ function and we are compiling for
82 /// the device, creates a diagnostic which is emitted if and when we realize
83 /// that the function will be codegen'ed.
84 ///
85 /// Example usage:
86 ///
87 /// // Variable-length arrays are not allowed in CUDA device code.
88 /// if (DiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentTarget())
89 /// return ExprError();
90 /// // Otherwise, continue parsing as normal.
92
93 /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
94 /// context is "used as host code".
95 ///
96 /// Same as DiagIfDeviceCode, with "host" and "device" switched.
98
99 /// Determines whether the given function is a CUDA device/host/kernel/etc.
100 /// function.
101 ///
102 /// Use this rather than examining the function's attributes yourself -- you
103 /// will get it wrong. Returns CUDAFunctionTarget::Host if D is null.
105 bool IgnoreImplicitHDAttr = false);
107
109 CVT_Device, /// Emitted on device side with a shadow variable on host side
110 CVT_Host, /// Emitted on host side only
111 CVT_Both, /// Emitted on both sides with different addresses
112 CVT_Unified, /// Emitted as a unified address, e.g. managed variables
113 };
114 /// Determines whether the given variable is emitted on host or device side.
116
117 /// Defines kinds of CUDA global host/device context where a function may be
118 /// called.
120 CTCK_Unknown, /// Unknown context
121 CTCK_InitGlobalVar, /// Function called during global variable
122 /// initialization
123 };
124
125 /// Define the current global CUDA host/device context where a function may be
126 /// called. Only used when a function is called outside of any functions.
130 Decl *D = nullptr;
132
137 Decl *D);
139 };
140
141 /// Gets the CUDA target for the current context.
143 return IdentifyTarget(dyn_cast<FunctionDecl>(SemaRef.CurContext));
144 }
145
146 static bool isImplicitHostDeviceFunction(const FunctionDecl *D);
147
148 // CUDA function call preference. Must be ordered numerically from
149 // worst to best.
151 CFP_Never, // Invalid caller/callee combination.
152 CFP_WrongSide, // Calls from host-device to host or device
153 // function that do not match current compilation
154 // mode.
155 CFP_HostDevice, // Any calls to host/device functions.
156 CFP_SameSide, // Calls from host-device to host or device
157 // function matching current compilation mode.
158 CFP_Native, // host-to-host or device-to-device calls.
159 };
160
161 /// Identifies relative preference of a given Caller/Callee
162 /// combination, based on their host/device attributes.
163 /// \param Caller function which needs address of \p Callee.
164 /// nullptr in case of global context.
165 /// \param Callee target function
166 ///
167 /// \returns preference value for particular Caller/Callee combination.
169 const FunctionDecl *Callee);
170
171 /// Determines whether Caller may invoke Callee, based on their CUDA
172 /// host/device attributes. Returns false if the call is not allowed.
173 ///
174 /// Note: Will return true for CFP_WrongSide calls. These may appear in
175 /// semantically correct CUDA programs, but only if they're never codegen'ed.
176 bool IsAllowedCall(const FunctionDecl *Caller, const FunctionDecl *Callee) {
177 return IdentifyPreference(Caller, Callee) != CFP_Never;
178 }
179
180 /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
181 /// depending on FD and the current compilation settings.
183
184 /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
185 /// and current compilation settings.
187
188 /// Check whether we're allowed to call Callee from the current context.
189 ///
190 /// - If the call is never allowed in a semantically-correct program
191 /// (CFP_Never), emits an error and returns false.
192 ///
193 /// - If the call is allowed in semantically-correct programs, but only if
194 /// it's never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to
195 /// be emitted if and when the caller is codegen'ed, and returns true.
196 ///
197 /// Will only create deferred diagnostics for a given SourceLocation once,
198 /// so you can safely call this multiple times without generating duplicate
199 /// deferred errors.
200 ///
201 /// - Otherwise, returns true without emitting any diagnostics.
203
205
206 /// Set __device__ or __host__ __device__ attributes on the given lambda
207 /// operator() method.
208 ///
209 /// CUDA lambdas by default is host device function unless it has explicit
210 /// host or device attribute.
211 void SetLambdaAttrs(CXXMethodDecl *Method);
212
213 /// Record \p FD if it is a CUDA/HIP implicit host device function used on
214 /// device side in device compilation.
216
217 /// Finds a function in \p Matches with highest calling priority
218 /// from \p Caller context and erases all functions with lower
219 /// calling priority.
221 const FunctionDecl *Caller,
222 llvm::SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>>
223 &Matches);
224
225 /// Given a implicit special member, infer its CUDA target from the
226 /// calls it needs to make to underlying base/field special members.
227 /// \param ClassDecl the class for which the member is being created.
228 /// \param CSM the kind of special member.
229 /// \param MemberDecl the special member itself.
230 /// \param ConstRHS true if this is a copy operation with a const object on
231 /// its RHS.
232 /// \param Diagnose true if this call should emit diagnostics.
233 /// \return true if there was an error inferring.
234 /// The result of this call is implicit CUDA target attribute(s) attached to
235 /// the member declaration.
238 CXXMethodDecl *MemberDecl,
239 bool ConstRHS, bool Diagnose);
240
241 /// \return true if \p CD can be considered empty according to CUDA
242 /// (E.2.3.1 in CUDA 7.5 Programming guide).
245
246 // \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
247 // case of error emits appropriate diagnostic and invalidates \p Var.
248 //
249 // \details CUDA allows only empty constructors as initializers for global
250 // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
251 // __shared__ variables whether they are local or not (they all are implicitly
252 // static in CUDA). One exception is that CUDA allows constant initializers
253 // for __constant__ and __device__ variables.
255
256 /// Check whether NewFD is a valid overload for CUDA. Emits
257 /// diagnostics and invalidates NewFD if not.
259 /// Copies target attributes from the template TD to the function FD.
261
262 /// Returns the name of the launch configuration function. This is the name
263 /// of the function that will be called to configure kernel call, with the
264 /// parameters specified via <<<>>>.
265 std::string getConfigureFuncName() const;
266
267private:
268 unsigned ForceHostDeviceDepth = 0;
269
270 friend class ASTReader;
271 friend class ASTWriter;
272};
273
274} // namespace clang
275
276namespace llvm {
277// Hash a FunctionDeclAndLoc by looking at both its FunctionDecl and its
278// SourceLocation.
279template <> struct DenseMapInfo<clang::SemaCUDA::FunctionDeclAndLoc> {
282 DenseMapInfo<clang::CanonicalDeclPtr<const clang::FunctionDecl>>;
283
285 return {FDBaseInfo::getEmptyKey(), clang::SourceLocation()};
286 }
287
289 return {FDBaseInfo::getTombstoneKey(), clang::SourceLocation()};
290 }
291
292 static unsigned getHashValue(const FunctionDeclAndLoc &FDL) {
293 return hash_combine(FDBaseInfo::getHashValue(FDL.FD),
294 FDL.Loc.getHashValue());
295 }
296
297 static bool isEqual(const FunctionDeclAndLoc &LHS,
298 const FunctionDeclAndLoc &RHS) {
299 return LHS.FD == RHS.FD && LHS.Loc == RHS.Loc;
300 }
301};
302} // namespace llvm
303
304#endif // LLVM_CLANG_SEMA_SEMACUDA_H
const Decl * D
Defines the C++ Decl subclasses, other than those for templates (found in DeclTemplate....
llvm::MachO::Target Target
Definition: MachO.h:51
SourceLocation Loc
Definition: SemaObjC.cpp:758
Defines the clang::SourceLocation class and associated facilities.
StateNode * Previous
Reads an AST files chain containing the contents of a translation unit.
Definition: ASTReader.h:378
Writes an AST file containing the contents of a translation unit.
Definition: ASTWriter.h:89
Represents a C++ constructor within a class.
Definition: DeclCXX.h:2539
Represents a C++ destructor within a class.
Definition: DeclCXX.h:2803
Represents a static or instance method of a struct/union/class.
Definition: DeclCXX.h:2064
Represents a C++ struct/union/class.
Definition: DeclCXX.h:258
A wrapper class around a pointer that always points to its canonical declaration.
Definition: Redeclarable.h:350
Decl - This represents one declaration (or definition), e.g.
Definition: DeclBase.h:86
Represents a function declaration or definition.
Definition: Decl.h:1932
Declaration of a template function.
Definition: DeclTemplate.h:957
Represents the results of name lookup.
Definition: Lookup.h:46
Scope - A scope is a transient data structure that is used while parsing the program.
Definition: Scope.h:41
A generic diagnostic builder for errors which may or may not be deferred.
Definition: SemaBase.h:110
Sema & SemaRef
Definition: SemaBase.h:40
void PushForceHostDevice()
Increments our count of the number of times we've seen a pragma forcing functions to be host device.
Definition: SemaCUDA.cpp:41
void checkAllowedInitializer(VarDecl *VD)
Definition: SemaCUDA.cpp:658
bool IsAllowedCall(const FunctionDecl *Caller, const FunctionDecl *Callee)
Determines whether Caller may invoke Callee, based on their CUDA host/device attributes.
Definition: SemaCUDA.h:176
void RecordImplicitHostDeviceFuncUsedByDevice(const FunctionDecl *FD)
Record FD if it is a CUDA/HIP implicit host device function used on device side in device compilation...
Definition: SemaCUDA.cpp:706
std::string getConfigureFuncName() const
Returns the name of the launch configuration function.
Definition: SemaCUDA.cpp:1070
bool PopForceHostDevice()
Decrements our count of the number of times we've seen a pragma forcing functions to be host device.
Definition: SemaCUDA.cpp:46
CUDAFunctionTarget IdentifyTarget(const FunctionDecl *D, bool IgnoreImplicitHDAttr=false)
Determines whether the given function is a CUDA device/host/kernel/etc.
Definition: SemaCUDA.cpp:136
void maybeAddHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous)
May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD, depending on FD and the current co...
Definition: SemaCUDA.cpp:740
ExprResult ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc, MultiExprArg ExecConfig, SourceLocation GGGLoc)
Definition: SemaCUDA.cpp:54
bool isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD)
Definition: SemaCUDA.cpp:512
bool isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *CD)
Definition: SemaCUDA.cpp:550
void checkTargetOverload(FunctionDecl *NewFD, const LookupResult &Previous)
Check whether NewFD is a valid overload for CUDA.
Definition: SemaCUDA.cpp:1005
CUDAFunctionTarget CurrentTarget()
Gets the CUDA target for the current context.
Definition: SemaCUDA.h:142
SemaDiagnosticBuilder DiagIfHostCode(SourceLocation Loc, unsigned DiagID)
Creates a SemaDiagnosticBuilder that emits the diagnostic if the current context is "used as host cod...
Definition: SemaCUDA.cpp:851
bool inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, CXXSpecialMemberKind CSM, CXXMethodDecl *MemberDecl, bool ConstRHS, bool Diagnose)
Given a implicit special member, infer its CUDA target from the calls it needs to make to underlying ...
Definition: SemaCUDA.cpp:372
struct clang::SemaCUDA::CUDATargetContext CurCUDATargetCtx
CUDATargetContextKind
Defines kinds of CUDA global host/device context where a function may be called.
Definition: SemaCUDA.h:119
@ CTCK_InitGlobalVar
Unknown context.
Definition: SemaCUDA.h:121
SemaDiagnosticBuilder DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID)
Creates a SemaDiagnosticBuilder that emits the diagnostic if the current context is "used as device c...
Definition: SemaCUDA.cpp:819
llvm::DenseSet< FunctionDeclAndLoc > LocsWithCUDACallDiags
FunctionDecls and SourceLocations for which CheckCall has emitted a (maybe deferred) "bad call" diagn...
Definition: SemaCUDA.h:63
bool CheckCall(SourceLocation Loc, FunctionDecl *Callee)
Check whether we're allowed to call Callee from the current context.
Definition: SemaCUDA.cpp:882
void inheritTargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD)
Copies target attributes from the template TD to the function FD.
Definition: SemaCUDA.cpp:1062
static bool isImplicitHostDeviceFunction(const FunctionDecl *D)
Definition: SemaCUDA.cpp:314
void CheckLambdaCapture(CXXMethodDecl *D, const sema::Capture &Capture)
Definition: SemaCUDA.cpp:952
void MaybeAddConstantAttr(VarDecl *VD)
May add implicit CUDAConstantAttr attribute to VD, depending on VD and current compilation settings.
Definition: SemaCUDA.cpp:805
void EraseUnwantedMatches(const FunctionDecl *Caller, llvm::SmallVectorImpl< std::pair< DeclAccessPair, FunctionDecl * > > &Matches)
Finds a function in Matches with highest calling priority from Caller context and erases all function...
Definition: SemaCUDA.cpp:320
llvm::DenseMap< CanonicalDeclPtr< const FunctionDecl >, FunctionDeclAndLoc > DeviceKnownEmittedFns
An inverse call graph, mapping known-emitted functions to one of their known-emitted callers (plus th...
Definition: SemaCUDA.h:72
void SetLambdaAttrs(CXXMethodDecl *Method)
Set device or host device attributes on the given lambda operator() method.
Definition: SemaCUDA.cpp:997
CUDAFunctionPreference IdentifyPreference(const FunctionDecl *Caller, const FunctionDecl *Callee)
Identifies relative preference of a given Caller/Callee combination, based on their host/device attri...
Definition: SemaCUDA.cpp:227
@ CVT_Host
Emitted on device side with a shadow variable on host side.
Definition: SemaCUDA.h:110
@ CVT_Both
Emitted on host side only.
Definition: SemaCUDA.h:111
@ CVT_Unified
Emitted on both sides with different addresses.
Definition: SemaCUDA.h:112
Sema - This implements semantic analysis and AST building for C.
Definition: Sema.h:535
DeclContext * CurContext
CurContext - This is the current declaration context of parsing.
Definition: Sema.h:1137
Encodes a location in the source.
unsigned getHashValue() const
Represents a variable declaration or definition.
Definition: Decl.h:879
The JSON file list parser is used to communicate input to InstallAPI.
CUDAFunctionTarget
Definition: Cuda.h:140
CXXSpecialMemberKind
Kinds of C++ special members.
Definition: Sema.h:445
Diagnostic wrappers for TextAPI types for error reporting.
Definition: Dominators.h:30
SemaCUDA::CUDATargetContext SavedCtx
Definition: SemaCUDA.h:135
Define the current global CUDA host/device context where a function may be called.
Definition: SemaCUDA.h:127
CUDATargetContextKind Kind
Definition: SemaCUDA.h:129
A pair of a canonical FunctionDecl and a SourceLocation.
Definition: SemaCUDA.h:55
CanonicalDeclPtr< const FunctionDecl > FD
Definition: SemaCUDA.h:56
DenseMapInfo< clang::CanonicalDeclPtr< const clang::FunctionDecl > > FDBaseInfo
Definition: SemaCUDA.h:282
static unsigned getHashValue(const FunctionDeclAndLoc &FDL)
Definition: SemaCUDA.h:292
static bool isEqual(const FunctionDeclAndLoc &LHS, const FunctionDeclAndLoc &RHS)
Definition: SemaCUDA.h:297