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
21 namespace clang {
22
SemaAMDGPU(Sema & S)23 SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
24
CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,CallExpr * TheCall)25 bool 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(SizeIdx);
34 [[maybe_unused]] ExprResult R =
35 SemaRef.VerifyIntegerConstantExpression(ArgExpr, &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(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
111 static bool
checkAMDGPUFlatWorkGroupSizeArguments(Sema & S,Expr * MinExpr,Expr * MaxExpr,const AMDGPUFlatWorkGroupSizeAttr & Attr)112 checkAMDGPUFlatWorkGroupSizeArguments(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(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
141 AMDGPUFlatWorkGroupSizeAttr *
CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)142 SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
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
addAMDGPUFlatWorkGroupSizeAttr(Decl * D,const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)153 void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
154 const AttributeCommonInfo &CI,
155 Expr *MinExpr, Expr *MaxExpr) {
156 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
157 D->addAttr(Attr);
158 }
159
handleAMDGPUFlatWorkGroupSizeAttr(Decl * D,const ParsedAttr & AL)160 void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
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
checkAMDGPUWavesPerEUArguments(Sema & S,Expr * MinExpr,Expr * MaxExpr,const AMDGPUWavesPerEUAttr & Attr)168 static 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
202 AMDGPUWavesPerEUAttr *
CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)203 SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
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
addAMDGPUWavesPerEUAttr(Decl * D,const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)214 void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
215 Expr *MinExpr, Expr *MaxExpr) {
216 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
217 D->addAttr(Attr);
218 }
219
handleAMDGPUWavesPerEUAttr(Decl * D,const ParsedAttr & AL)220 void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
221 if (!AL.checkAtLeastNumArgs(SemaRef, 1) || !AL.checkAtMostNumArgs(SemaRef, 2))
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
handleAMDGPUNumSGPRAttr(Decl * D,const ParsedAttr & AL)230 void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
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
handleAMDGPUNumVGPRAttr(Decl * D,const ParsedAttr & AL)240 void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
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
250 static bool
checkAMDGPUMaxNumWorkGroupsArguments(Sema & S,Expr * XExpr,Expr * YExpr,Expr * ZExpr,const AMDGPUMaxNumWorkGroupsAttr & Attr)251 checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
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
CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo & CI,Expr * XExpr,Expr * YExpr,Expr * ZExpr)283 AMDGPUMaxNumWorkGroupsAttr *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(SemaRef, XExpr, YExpr, ZExpr,
289 TmpAttr))
290 return nullptr;
291
292 return ::new (Context)
293 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
294 }
295
addAMDGPUMaxNumWorkGroupsAttr(Decl * D,const AttributeCommonInfo & CI,Expr * XExpr,Expr * YExpr,Expr * ZExpr)296 void 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(Attr);
302 }
303
handleAMDGPUMaxNumWorkGroupsAttr(Decl * D,const ParsedAttr & AL)304 void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
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
312