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
30 const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
31 assert(FD && "AMDGPU builtins should not be used outside of a function");
32 llvm::StringMap<bool> CallerFeatureMap;
33 getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD);
34 bool HasGFX950Insts =
35 Builtin::evaluateRequiredTargetFeatures("gfx950-insts", CallerFeatureMap);
36
37 switch (BuiltinID) {
38 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_lds:
39 case AMDGPU::BI__builtin_amdgcn_load_to_lds:
40 case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
41 constexpr const int SizeIdx = 2;
42 llvm::APSInt Size;
43 Expr *ArgExpr = TheCall->getArg(SizeIdx);
44 [[maybe_unused]] ExprResult R =
45 SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
46 assert(!R.isInvalid());
47 switch (Size.getSExtValue()) {
48 case 1:
49 case 2:
50 case 4:
51 return false;
52 case 12:
53 case 16: {
54 if (HasGFX950Insts)
55 return false;
56 [[fallthrough]];
57 }
58 default:
59 Diag(ArgExpr->getExprLoc(), diag::err_amdgcn_load_lds_size_invalid_value)
60 << ArgExpr->getSourceRange();
61 Diag(ArgExpr->getExprLoc(), diag::note_amdgcn_load_lds_size_valid_value)
62 << HasGFX950Insts << ArgExpr->getSourceRange();
63 return true;
64 }
65 }
66 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
67 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
68 return false;
69 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
70 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
71 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
72 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
73 OrderIndex = 2;
74 ScopeIndex = 3;
75 break;
76 case AMDGPU::BI__builtin_amdgcn_fence:
77 OrderIndex = 0;
78 ScopeIndex = 1;
79 break;
80 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
81 return checkMovDPPFunctionCall(TheCall, 5, 1);
82 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
83 return checkMovDPPFunctionCall(TheCall, 2, 1);
84 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
85 return checkMovDPPFunctionCall(TheCall, 6, 2);
86 }
87 default:
88 return false;
89 }
90
91 ExprResult Arg = TheCall->getArg(OrderIndex);
92 auto ArgExpr = Arg.get();
93 Expr::EvalResult ArgResult;
94
95 if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))
96 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
97 << ArgExpr->getType();
98 auto Ord = ArgResult.Val.getInt().getZExtValue();
99
100 // Check validity of memory ordering as per C11 / C++11's memody model.
101 // Only fence needs check. Atomic dec/inc allow all memory orders.
102 if (!llvm::isValidAtomicOrderingCABI(Ord))
103 return Diag(ArgExpr->getBeginLoc(),
104 diag::warn_atomic_op_has_invalid_memory_order)
105 << 0 << ArgExpr->getSourceRange();
106 switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
107 case llvm::AtomicOrderingCABI::relaxed:
108 case llvm::AtomicOrderingCABI::consume:
109 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
110 return Diag(ArgExpr->getBeginLoc(),
111 diag::warn_atomic_op_has_invalid_memory_order)
112 << 0 << ArgExpr->getSourceRange();
113 break;
114 case llvm::AtomicOrderingCABI::acquire:
115 case llvm::AtomicOrderingCABI::release:
116 case llvm::AtomicOrderingCABI::acq_rel:
117 case llvm::AtomicOrderingCABI::seq_cst:
118 break;
119 }
120
121 Arg = TheCall->getArg(ScopeIndex);
122 ArgExpr = Arg.get();
123 Expr::EvalResult ArgResult1;
124 // Check that sync scope is a constant literal
125 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))
126 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
127 << ArgExpr->getType();
128
129 return false;
130 }
131
checkMovDPPFunctionCall(CallExpr * TheCall,unsigned NumArgs,unsigned NumDataArgs)132 bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
133 unsigned NumDataArgs) {
134 assert(NumDataArgs <= 2);
135 if (SemaRef.checkArgCountRange(TheCall, NumArgs, NumArgs))
136 return true;
137 Expr *Args[2];
138 QualType ArgTys[2];
139 for (unsigned I = 0; I != NumDataArgs; ++I) {
140 Args[I] = TheCall->getArg(I);
141 ArgTys[I] = Args[I]->getType();
142 // TODO: Vectors can also be supported.
143 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
144 SemaRef.Diag(Args[I]->getBeginLoc(),
145 diag::err_typecheck_cond_expect_int_float)
146 << ArgTys[I] << Args[I]->getSourceRange();
147 return true;
148 }
149 }
150 if (NumDataArgs < 2)
151 return false;
152
153 if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1]))
154 return false;
155
156 if (((ArgTys[0]->isUnsignedIntegerType() &&
157 ArgTys[1]->isSignedIntegerType()) ||
158 (ArgTys[0]->isSignedIntegerType() &&
159 ArgTys[1]->isUnsignedIntegerType())) &&
160 getASTContext().getTypeSize(ArgTys[0]) ==
161 getASTContext().getTypeSize(ArgTys[1]))
162 return false;
163
164 SemaRef.Diag(Args[1]->getBeginLoc(),
165 diag::err_typecheck_call_different_arg_types)
166 << ArgTys[0] << ArgTys[1];
167 return true;
168 }
169
170 static bool
checkAMDGPUFlatWorkGroupSizeArguments(Sema & S,Expr * MinExpr,Expr * MaxExpr,const AMDGPUFlatWorkGroupSizeAttr & Attr)171 checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
172 const AMDGPUFlatWorkGroupSizeAttr &Attr) {
173 // Accept template arguments for now as they depend on something else.
174 // We'll get to check them when they eventually get instantiated.
175 if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
176 return false;
177
178 uint32_t Min = 0;
179 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
180 return true;
181
182 uint32_t Max = 0;
183 if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
184 return true;
185
186 if (Min == 0 && Max != 0) {
187 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
188 << &Attr << 0;
189 return true;
190 }
191 if (Min > Max) {
192 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
193 << &Attr << 1;
194 return true;
195 }
196
197 return false;
198 }
199
200 AMDGPUFlatWorkGroupSizeAttr *
CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)201 SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
202 Expr *MinExpr, Expr *MaxExpr) {
203 ASTContext &Context = getASTContext();
204 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
205
206 if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
207 return nullptr;
208 return ::new (Context)
209 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
210 }
211
addAMDGPUFlatWorkGroupSizeAttr(Decl * D,const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)212 void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
213 const AttributeCommonInfo &CI,
214 Expr *MinExpr, Expr *MaxExpr) {
215 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
216 D->addAttr(Attr);
217 }
218
handleAMDGPUFlatWorkGroupSizeAttr(Decl * D,const ParsedAttr & AL)219 void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
220 const ParsedAttr &AL) {
221 Expr *MinExpr = AL.getArgAsExpr(0);
222 Expr *MaxExpr = AL.getArgAsExpr(1);
223
224 addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
225 }
226
checkAMDGPUWavesPerEUArguments(Sema & S,Expr * MinExpr,Expr * MaxExpr,const AMDGPUWavesPerEUAttr & Attr)227 static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
228 Expr *MaxExpr,
229 const AMDGPUWavesPerEUAttr &Attr) {
230 if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
231 (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
232 return true;
233
234 // Accept template arguments for now as they depend on something else.
235 // We'll get to check them when they eventually get instantiated.
236 if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
237 return false;
238
239 uint32_t Min = 0;
240 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
241 return true;
242
243 uint32_t Max = 0;
244 if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
245 return true;
246
247 if (Min == 0 && Max != 0) {
248 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
249 << &Attr << 0;
250 return true;
251 }
252 if (Max != 0 && Min > Max) {
253 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
254 << &Attr << 1;
255 return true;
256 }
257
258 return false;
259 }
260
261 AMDGPUWavesPerEUAttr *
CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)262 SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
263 Expr *MinExpr, Expr *MaxExpr) {
264 ASTContext &Context = getASTContext();
265 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
266
267 if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
268 return nullptr;
269
270 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
271 }
272
addAMDGPUWavesPerEUAttr(Decl * D,const AttributeCommonInfo & CI,Expr * MinExpr,Expr * MaxExpr)273 void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
274 Expr *MinExpr, Expr *MaxExpr) {
275 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
276 D->addAttr(Attr);
277 }
278
handleAMDGPUWavesPerEUAttr(Decl * D,const ParsedAttr & AL)279 void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
280 if (!AL.checkAtLeastNumArgs(SemaRef, 1) || !AL.checkAtMostNumArgs(SemaRef, 2))
281 return;
282
283 Expr *MinExpr = AL.getArgAsExpr(0);
284 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
285
286 addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
287 }
288
handleAMDGPUNumSGPRAttr(Decl * D,const ParsedAttr & AL)289 void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
290 uint32_t NumSGPR = 0;
291 Expr *NumSGPRExpr = AL.getArgAsExpr(0);
292 if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
293 return;
294
295 D->addAttr(::new (getASTContext())
296 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
297 }
298
handleAMDGPUNumVGPRAttr(Decl * D,const ParsedAttr & AL)299 void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
300 uint32_t NumVGPR = 0;
301 Expr *NumVGPRExpr = AL.getArgAsExpr(0);
302 if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
303 return;
304
305 D->addAttr(::new (getASTContext())
306 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
307 }
308
309 static bool
checkAMDGPUMaxNumWorkGroupsArguments(Sema & S,Expr * XExpr,Expr * YExpr,Expr * ZExpr,const AMDGPUMaxNumWorkGroupsAttr & Attr)310 checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
311 Expr *ZExpr,
312 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
313 if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
314 (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
315 (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
316 return true;
317
318 // Accept template arguments for now as they depend on something else.
319 // We'll get to check them when they eventually get instantiated.
320 if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
321 (ZExpr && ZExpr->isValueDependent()))
322 return false;
323
324 uint32_t NumWG = 0;
325 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
326 for (int i = 0; i < 3; i++) {
327 if (Exprs[i]) {
328 if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
329 /*StrictlyUnsigned=*/true))
330 return true;
331 if (NumWG == 0) {
332 S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
333 << &Attr << Exprs[i]->getSourceRange();
334 return true;
335 }
336 }
337 }
338
339 return false;
340 }
341
CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo & CI,Expr * XExpr,Expr * YExpr,Expr * ZExpr)342 AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
343 const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
344 ASTContext &Context = getASTContext();
345 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
346
347 if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
348 TmpAttr))
349 return nullptr;
350
351 return ::new (Context)
352 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
353 }
354
addAMDGPUMaxNumWorkGroupsAttr(Decl * D,const AttributeCommonInfo & CI,Expr * XExpr,Expr * YExpr,Expr * ZExpr)355 void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
356 const AttributeCommonInfo &CI,
357 Expr *XExpr, Expr *YExpr,
358 Expr *ZExpr) {
359 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
360 D->addAttr(Attr);
361 }
362
handleAMDGPUMaxNumWorkGroupsAttr(Decl * D,const ParsedAttr & AL)363 void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
364 const ParsedAttr &AL) {
365 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
366 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
367 addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
368 }
369
370 } // namespace clang
371