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
13#include "clang/Sema/SemaAMDGPU.h"
14#include "clang/Basic/DiagnosticSema.h"
15#include "clang/Basic/TargetBuiltins.h"
16#include "clang/Sema/Ownership.h"
17#include "clang/Sema/Sema.h"
18#include "llvm/Support/AtomicOrdering.h"
19#include <cstdint>
20
21namespace clang {
22
23SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
24
25bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
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(Arg: SizeIdx);
34 [[maybe_unused]] ExprResult R =
35 SemaRef.VerifyIntegerConstantExpression(E: ArgExpr, Result: &Size);
36 assert(!R.isInvalid());
37 switch (Size.getSExtValue()) {
38 case 1:
39 case 2:
40 case 4:
41 return false;
42 default:
43 Diag(Loc: ArgExpr->getExprLoc(),
44 DiagID: diag::err_amdgcn_global_load_lds_size_invalid_value)
45 << ArgExpr->getSourceRange();
46 Diag(Loc: ArgExpr->getExprLoc(),
47 DiagID: 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(Arg: OrderIndex);
71 auto ArgExpr = Arg.get();
72 Expr::EvalResult ArgResult;
73
74 if (!ArgExpr->EvaluateAsInt(Result&: ArgResult, Ctx: getASTContext()))
75 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: 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(I: Ord))
82 return Diag(Loc: ArgExpr->getBeginLoc(),
83 DiagID: 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(Loc: ArgExpr->getBeginLoc(),
90 DiagID: 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(Arg: ScopeIndex);
101 ArgExpr = Arg.get();
102 Expr::EvalResult ArgResult1;
103 // Check that sync scope is a constant literal
104 if (!ArgExpr->EvaluateAsConstantExpr(Result&: ArgResult1, Ctx: getASTContext()))
105 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: diag::err_expr_not_string_literal)
106 << ArgExpr->getType();
107
108 return false;
109}
110
111static bool
112checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
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(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
121 return true;
122
123 uint32_t Max = 0;
124 if (!S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
125 return true;
126
127 if (Min == 0 && Max != 0) {
128 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
129 << &Attr << 0;
130 return true;
131 }
132 if (Min > Max) {
133 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
134 << &Attr << 1;
135 return true;
136 }
137
138 return false;
139}
140
141AMDGPUFlatWorkGroupSizeAttr *
142SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
143 Expr *MinExpr, Expr *MaxExpr) {
144 ASTContext &Context = getASTContext();
145 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
146
147 if (checkAMDGPUFlatWorkGroupSizeArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
148 return nullptr;
149 return ::new (Context)
150 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
151}
152
153void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
154 const AttributeCommonInfo &CI,
155 Expr *MinExpr, Expr *MaxExpr) {
156 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
157 D->addAttr(A: Attr);
158}
159
160void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
161 const ParsedAttr &AL) {
162 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
163 Expr *MaxExpr = AL.getArgAsExpr(Arg: 1);
164
165 addAMDGPUFlatWorkGroupSizeAttr(D, CI: AL, MinExpr, MaxExpr);
166}
167
168static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
169 Expr *MaxExpr,
170 const AMDGPUWavesPerEUAttr &Attr) {
171 if (S.DiagnoseUnexpandedParameterPack(E: MinExpr) ||
172 (MaxExpr && S.DiagnoseUnexpandedParameterPack(E: 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(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
182 return true;
183
184 uint32_t Max = 0;
185 if (MaxExpr && !S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
186 return true;
187
188 if (Min == 0 && Max != 0) {
189 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
190 << &Attr << 0;
191 return true;
192 }
193 if (Max != 0 && Min > Max) {
194 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
195 << &Attr << 1;
196 return true;
197 }
198
199 return false;
200}
201
202AMDGPUWavesPerEUAttr *
203SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
204 Expr *MinExpr, Expr *MaxExpr) {
205 ASTContext &Context = getASTContext();
206 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
207
208 if (checkAMDGPUWavesPerEUArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
209 return nullptr;
210
211 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
212}
213
214void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
215 Expr *MinExpr, Expr *MaxExpr) {
216 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
217 D->addAttr(A: Attr);
218}
219
220void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
221 if (!AL.checkAtLeastNumArgs(S&: SemaRef, Num: 1) || !AL.checkAtMostNumArgs(S&: SemaRef, Num: 2))
222 return;
223
224 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
225 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
226
227 addAMDGPUWavesPerEUAttr(D, CI: AL, MinExpr, MaxExpr);
228}
229
230void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
231 uint32_t NumSGPR = 0;
232 Expr *NumSGPRExpr = AL.getArgAsExpr(Arg: 0);
233 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumSGPRExpr, Val&: NumSGPR))
234 return;
235
236 D->addAttr(A: ::new (getASTContext())
237 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
238}
239
240void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
241 uint32_t NumVGPR = 0;
242 Expr *NumVGPRExpr = AL.getArgAsExpr(Arg: 0);
243 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumVGPRExpr, Val&: NumVGPR))
244 return;
245
246 D->addAttr(A: ::new (getASTContext())
247 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
248}
249
250static bool
251checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
252 Expr *ZExpr,
253 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
254 if (S.DiagnoseUnexpandedParameterPack(E: XExpr) ||
255 (YExpr && S.DiagnoseUnexpandedParameterPack(E: YExpr)) ||
256 (ZExpr && S.DiagnoseUnexpandedParameterPack(E: 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(AI: Attr, Expr: Exprs[i], Val&: NumWG, Idx: i,
270 /*StrictlyUnsigned=*/true))
271 return true;
272 if (NumWG == 0) {
273 S.Diag(Loc: Attr.getLoc(), DiagID: diag::err_attribute_argument_is_zero)
274 << &Attr << Exprs[i]->getSourceRange();
275 return true;
276 }
277 }
278 }
279
280 return false;
281}
282
283AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
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(S&: SemaRef, XExpr, YExpr, ZExpr,
289 Attr: TmpAttr))
290 return nullptr;
291
292 return ::new (Context)
293 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
294}
295
296void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
297 const AttributeCommonInfo &CI,
298 Expr *XExpr, Expr *YExpr,
299 Expr *ZExpr) {
300 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
301 D->addAttr(A: Attr);
302}
303
304void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
305 const ParsedAttr &AL) {
306 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
307 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(Arg: 2) : nullptr;
308 addAMDGPUMaxNumWorkGroupsAttr(D, CI: AL, XExpr: AL.getArgAsExpr(Arg: 0), YExpr, ZExpr);
309}
310
311} // namespace clang
312