clang 20.0.0git
SemaAMDGPU.cpp
Go to the documentation of this file.
1//===------ SemaAMDGPU.cpp ------- AMDGPU target-specific routines --------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file implements semantic analysis functions specific to AMDGPU.
10//
11//===----------------------------------------------------------------------===//
12
17#include "clang/Sema/Sema.h"
18#include "llvm/Support/AtomicOrdering.h"
19#include <cstdint>
20
21namespace clang {
22
24
26 CallExpr *TheCall) {
27 // position of memory order and scope arguments in the builtin
28 unsigned OrderIndex, ScopeIndex;
29 switch (BuiltinID) {
30 case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
31 constexpr const int SizeIdx = 2;
32 llvm::APSInt Size;
33 Expr *ArgExpr = TheCall->getArg(SizeIdx);
34 [[maybe_unused]] ExprResult R =
36 assert(!R.isInvalid());
37 switch (Size.getSExtValue()) {
38 case 1:
39 case 2:
40 case 4:
41 return false;
42 default:
43 Diag(ArgExpr->getExprLoc(),
44 diag::err_amdgcn_global_load_lds_size_invalid_value)
45 << ArgExpr->getSourceRange();
46 Diag(ArgExpr->getExprLoc(),
47 diag::note_amdgcn_global_load_lds_size_valid_value)
48 << ArgExpr->getSourceRange();
49 return true;
50 }
51 }
52 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
53 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
54 return false;
55 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
56 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
57 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
58 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
59 OrderIndex = 2;
60 ScopeIndex = 3;
61 break;
62 case AMDGPU::BI__builtin_amdgcn_fence:
63 OrderIndex = 0;
64 ScopeIndex = 1;
65 break;
66 default:
67 return false;
68 }
69
70 ExprResult Arg = TheCall->getArg(OrderIndex);
71 auto ArgExpr = Arg.get();
72 Expr::EvalResult ArgResult;
73
74 if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))
75 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
76 << ArgExpr->getType();
77 auto Ord = ArgResult.Val.getInt().getZExtValue();
78
79 // Check validity of memory ordering as per C11 / C++11's memody model.
80 // Only fence needs check. Atomic dec/inc allow all memory orders.
81 if (!llvm::isValidAtomicOrderingCABI(Ord))
82 return Diag(ArgExpr->getBeginLoc(),
83 diag::warn_atomic_op_has_invalid_memory_order)
84 << 0 << ArgExpr->getSourceRange();
85 switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
86 case llvm::AtomicOrderingCABI::relaxed:
87 case llvm::AtomicOrderingCABI::consume:
88 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
89 return Diag(ArgExpr->getBeginLoc(),
90 diag::warn_atomic_op_has_invalid_memory_order)
91 << 0 << ArgExpr->getSourceRange();
92 break;
93 case llvm::AtomicOrderingCABI::acquire:
94 case llvm::AtomicOrderingCABI::release:
95 case llvm::AtomicOrderingCABI::acq_rel:
96 case llvm::AtomicOrderingCABI::seq_cst:
97 break;
98 }
99
100 Arg = TheCall->getArg(ScopeIndex);
101 ArgExpr = Arg.get();
102 Expr::EvalResult ArgResult1;
103 // Check that sync scope is a constant literal
104 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))
105 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
106 << ArgExpr->getType();
107
108 return false;
109}
110
111static bool
113 const AMDGPUFlatWorkGroupSizeAttr &Attr) {
114 // Accept template arguments for now as they depend on something else.
115 // We'll get to check them when they eventually get instantiated.
116 if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
117 return false;
118
119 uint32_t Min = 0;
120 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
121 return true;
122
123 uint32_t Max = 0;
124 if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
125 return true;
126
127 if (Min == 0 && Max != 0) {
128 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
129 << &Attr << 0;
130 return true;
131 }
132 if (Min > Max) {
133 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
134 << &Attr << 1;
135 return true;
136 }
137
138 return false;
139}
140
141AMDGPUFlatWorkGroupSizeAttr *
143 Expr *MinExpr, Expr *MaxExpr) {
144 ASTContext &Context = getASTContext();
145 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
146
147 if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
148 return nullptr;
149 return ::new (Context)
150 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
151}
152
154 const AttributeCommonInfo &CI,
155 Expr *MinExpr, Expr *MaxExpr) {
156 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
157 D->addAttr(Attr);
158}
159
161 const ParsedAttr &AL) {
162 Expr *MinExpr = AL.getArgAsExpr(0);
163 Expr *MaxExpr = AL.getArgAsExpr(1);
164
165 addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
166}
167
168static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
169 Expr *MaxExpr,
170 const AMDGPUWavesPerEUAttr &Attr) {
171 if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
172 (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
173 return true;
174
175 // Accept template arguments for now as they depend on something else.
176 // We'll get to check them when they eventually get instantiated.
177 if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
178 return false;
179
180 uint32_t Min = 0;
181 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
182 return true;
183
184 uint32_t Max = 0;
185 if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
186 return true;
187
188 if (Min == 0 && Max != 0) {
189 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
190 << &Attr << 0;
191 return true;
192 }
193 if (Max != 0 && Min > Max) {
194 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
195 << &Attr << 1;
196 return true;
197 }
198
199 return false;
200}
201
202AMDGPUWavesPerEUAttr *
204 Expr *MinExpr, Expr *MaxExpr) {
205 ASTContext &Context = getASTContext();
206 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
207
208 if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
209 return nullptr;
210
211 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
212}
213
215 Expr *MinExpr, Expr *MaxExpr) {
216 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
217 D->addAttr(Attr);
218}
219
222 return;
223
224 Expr *MinExpr = AL.getArgAsExpr(0);
225 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
226
227 addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
228}
229
231 uint32_t NumSGPR = 0;
232 Expr *NumSGPRExpr = AL.getArgAsExpr(0);
233 if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
234 return;
235
236 D->addAttr(::new (getASTContext())
237 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
238}
239
241 uint32_t NumVGPR = 0;
242 Expr *NumVGPRExpr = AL.getArgAsExpr(0);
243 if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
244 return;
245
246 D->addAttr(::new (getASTContext())
247 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
248}
249
250static bool
252 Expr *ZExpr,
253 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
254 if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
255 (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
256 (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
257 return true;
258
259 // Accept template arguments for now as they depend on something else.
260 // We'll get to check them when they eventually get instantiated.
261 if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
262 (ZExpr && ZExpr->isValueDependent()))
263 return false;
264
265 uint32_t NumWG = 0;
266 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
267 for (int i = 0; i < 3; i++) {
268 if (Exprs[i]) {
269 if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
270 /*StrictlyUnsigned=*/true))
271 return true;
272 if (NumWG == 0) {
273 S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
274 << &Attr << Exprs[i]->getSourceRange();
275 return true;
276 }
277 }
278 }
279
280 return false;
281}
282
284 const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
285 ASTContext &Context = getASTContext();
286 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
287
288 if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
289 TmpAttr))
290 return nullptr;
291
292 return ::new (Context)
293 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
294}
295
297 const AttributeCommonInfo &CI,
298 Expr *XExpr, Expr *YExpr,
299 Expr *ZExpr) {
300 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
301 D->addAttr(Attr);
302}
303
305 const ParsedAttr &AL) {
306 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
307 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
308 addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
309}
310
311} // namespace clang
const Decl * D
This file declares semantic analysis functions specific to AMDGPU.
Enumerates target-specific builtins in their own namespaces within namespace clang.
APSInt & getInt()
Definition: APValue.h:423
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
Definition: ASTContext.h:187
PtrTy get() const
Definition: Ownership.h:170
bool isInvalid() const
Definition: Ownership.h:166
Attr - This represents one attribute.
Definition: Attr.h:42
SourceLocation getLocation() const
Definition: Attr.h:95
SourceLocation getLoc() const
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition: Expr.h:2830
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition: Expr.h:3021
Decl - This represents one declaration (or definition), e.g.
Definition: DeclBase.h:86
This represents one expression.
Definition: Expr.h:110
bool isValueDependent() const
Determines whether the value of this expression depends on.
Definition: Expr.h:175
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
Definition: Expr.cpp:277
ParsedAttr - Represents a syntactic attribute.
Definition: ParsedAttr.h:129
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
Definition: ParsedAttr.h:386
Expr * getArgAsExpr(unsigned Arg) const
Definition: ParsedAttr.h:398
bool checkAtLeastNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has at least as many args as Num.
Definition: ParsedAttr.cpp:303
bool checkAtMostNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has at most as many args as Num.
Definition: ParsedAttr.cpp:308
void handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, const ParsedAttr &AL)
Definition: SemaAMDGPU.cpp:304
void addAMDGPUFlatWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size attribute to a particular declar...
Definition: SemaAMDGPU.cpp:153
SemaAMDGPU(Sema &S)
Definition: SemaAMDGPU.cpp:23
void handleAMDGPUFlatWorkGroupSizeAttr(Decl *D, const ParsedAttr &AL)
Definition: SemaAMDGPU.cpp:160
void handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL)
Definition: SemaAMDGPU.cpp:230
AMDGPUMaxNumWorkGroupsAttr * CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr)
Create an AMDGPUMaxNumWorkGroupsAttr attribute.
Definition: SemaAMDGPU.cpp:283
AMDGPUWavesPerEUAttr * CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
Create an AMDGPUWavesPerEUAttr attribute.
Definition: SemaAMDGPU.cpp:203
void handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL)
Definition: SemaAMDGPU.cpp:240
AMDGPUFlatWorkGroupSizeAttr * CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
Create an AMDGPUWavesPerEUAttr attribute.
Definition: SemaAMDGPU.cpp:142
void handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL)
Definition: SemaAMDGPU.cpp:220
bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)
Definition: SemaAMDGPU.cpp:25
void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a particular declaration.
Definition: SemaAMDGPU.cpp:214
void addAMDGPUMaxNumWorkGroupsAttr(Decl *D, const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr)
addAMDGPUMaxNumWorkGroupsAttr - Adds an amdgpu_max_num_work_groups attribute to a particular declarat...
Definition: SemaAMDGPU.cpp:296
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID, bool DeferHint=false)
Emit a diagnostic.
Definition: SemaBase.cpp:60
ASTContext & getASTContext() const
Definition: SemaBase.cpp:9
Sema & SemaRef
Definition: SemaBase.h:40
Sema - This implements semantic analysis and AST building for C.
Definition: Sema.h:493
ExprResult VerifyIntegerConstantExpression(Expr *E, llvm::APSInt *Result, VerifyICEDiagnoser &Diagnoser, AllowFoldKind CanFold=NoFold)
VerifyIntegerConstantExpression - Verifies that an expression is an ICE, and reports the appropriate ...
Definition: SemaExpr.cpp:16933
bool DiagnoseUnexpandedParameterPack(SourceLocation Loc, TypeSourceInfo *T, UnexpandedParameterPackContext UPPC)
If the given type contains an unexpanded parameter pack, diagnose the error.
bool checkUInt32Argument(const AttrInfo &AI, const Expr *Expr, uint32_t &Val, unsigned Idx=UINT_MAX, bool StrictlyUnsigned=false)
If Expr is a valid integer constant, get the value of the integer expression and return success or fa...
Definition: Sema.h:4434
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
Definition: Stmt.cpp:326
The JSON file list parser is used to communicate input to InstallAPI.
static bool checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr, Expr *ZExpr, const AMDGPUMaxNumWorkGroupsAttr &Attr)
Definition: SemaAMDGPU.cpp:251
static bool checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUFlatWorkGroupSizeAttr &Attr)
Definition: SemaAMDGPU.cpp:112
static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUWavesPerEUAttr &Attr)
Definition: SemaAMDGPU.cpp:168
EvalResult is a struct with detailed info about an evaluated expression.
Definition: Expr.h:642
APValue Val
Val - This is the value the expression can be folded to.
Definition: Expr.h:644