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
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(FeatureMap&: CallerFeatureMap, FD);
34 bool HasGFX950Insts =
35 Builtin::evaluateRequiredTargetFeatures(RequiredFatures: "gfx950-insts", TargetFetureMap: 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(Arg: SizeIdx);
44 [[maybe_unused]] ExprResult R =
45 SemaRef.VerifyIntegerConstantExpression(E: ArgExpr, Result: &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(Loc: ArgExpr->getExprLoc(), DiagID: diag::err_amdgcn_load_lds_size_invalid_value)
60 << ArgExpr->getSourceRange();
61 Diag(Loc: ArgExpr->getExprLoc(), DiagID: 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, NumArgs: 5, NumDataArgs: 1);
82 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
83 return checkMovDPPFunctionCall(TheCall, NumArgs: 2, NumDataArgs: 1);
84 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
85 return checkMovDPPFunctionCall(TheCall, NumArgs: 6, NumDataArgs: 2);
86 }
87 default:
88 return false;
89 }
90
91 ExprResult Arg = TheCall->getArg(Arg: OrderIndex);
92 auto ArgExpr = Arg.get();
93 Expr::EvalResult ArgResult;
94
95 if (!ArgExpr->EvaluateAsInt(Result&: ArgResult, Ctx: getASTContext()))
96 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: 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(I: Ord))
103 return Diag(Loc: ArgExpr->getBeginLoc(),
104 DiagID: 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(Loc: ArgExpr->getBeginLoc(),
111 DiagID: 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(Arg: ScopeIndex);
122 ArgExpr = Arg.get();
123 Expr::EvalResult ArgResult1;
124 // Check that sync scope is a constant literal
125 if (!ArgExpr->EvaluateAsConstantExpr(Result&: ArgResult1, Ctx: getASTContext()))
126 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: diag::err_expr_not_string_literal)
127 << ArgExpr->getType();
128
129 return false;
130}
131
132bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
133 unsigned NumDataArgs) {
134 assert(NumDataArgs <= 2);
135 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: NumArgs, MaxArgCount: 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(Arg: 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(Loc: Args[I]->getBeginLoc(),
145 DiagID: 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(T1: ArgTys[0], T2: 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(T: ArgTys[0]) ==
161 getASTContext().getTypeSize(T: ArgTys[1]))
162 return false;
163
164 SemaRef.Diag(Loc: Args[1]->getBeginLoc(),
165 DiagID: diag::err_typecheck_call_different_arg_types)
166 << ArgTys[0] << ArgTys[1];
167 return true;
168}
169
170static bool
171checkAMDGPUFlatWorkGroupSizeArguments(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(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
180 return true;
181
182 uint32_t Max = 0;
183 if (!S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
184 return true;
185
186 if (Min == 0 && Max != 0) {
187 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
188 << &Attr << 0;
189 return true;
190 }
191 if (Min > Max) {
192 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
193 << &Attr << 1;
194 return true;
195 }
196
197 return false;
198}
199
200AMDGPUFlatWorkGroupSizeAttr *
201SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
202 Expr *MinExpr, Expr *MaxExpr) {
203 ASTContext &Context = getASTContext();
204 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
205
206 if (checkAMDGPUFlatWorkGroupSizeArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
207 return nullptr;
208 return ::new (Context)
209 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
210}
211
212void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
213 const AttributeCommonInfo &CI,
214 Expr *MinExpr, Expr *MaxExpr) {
215 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
216 D->addAttr(A: Attr);
217}
218
219void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
220 const ParsedAttr &AL) {
221 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
222 Expr *MaxExpr = AL.getArgAsExpr(Arg: 1);
223
224 addAMDGPUFlatWorkGroupSizeAttr(D, CI: AL, MinExpr, MaxExpr);
225}
226
227static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
228 Expr *MaxExpr,
229 const AMDGPUWavesPerEUAttr &Attr) {
230 if (S.DiagnoseUnexpandedParameterPack(E: MinExpr) ||
231 (MaxExpr && S.DiagnoseUnexpandedParameterPack(E: 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(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
241 return true;
242
243 uint32_t Max = 0;
244 if (MaxExpr && !S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
245 return true;
246
247 if (Min == 0 && Max != 0) {
248 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
249 << &Attr << 0;
250 return true;
251 }
252 if (Max != 0 && Min > Max) {
253 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
254 << &Attr << 1;
255 return true;
256 }
257
258 return false;
259}
260
261AMDGPUWavesPerEUAttr *
262SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
263 Expr *MinExpr, Expr *MaxExpr) {
264 ASTContext &Context = getASTContext();
265 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
266
267 if (checkAMDGPUWavesPerEUArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
268 return nullptr;
269
270 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
271}
272
273void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
274 Expr *MinExpr, Expr *MaxExpr) {
275 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
276 D->addAttr(A: Attr);
277}
278
279void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
280 if (!AL.checkAtLeastNumArgs(S&: SemaRef, Num: 1) || !AL.checkAtMostNumArgs(S&: SemaRef, Num: 2))
281 return;
282
283 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
284 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
285
286 addAMDGPUWavesPerEUAttr(D, CI: AL, MinExpr, MaxExpr);
287}
288
289void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
290 uint32_t NumSGPR = 0;
291 Expr *NumSGPRExpr = AL.getArgAsExpr(Arg: 0);
292 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumSGPRExpr, Val&: NumSGPR))
293 return;
294
295 D->addAttr(A: ::new (getASTContext())
296 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
297}
298
299void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
300 uint32_t NumVGPR = 0;
301 Expr *NumVGPRExpr = AL.getArgAsExpr(Arg: 0);
302 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumVGPRExpr, Val&: NumVGPR))
303 return;
304
305 D->addAttr(A: ::new (getASTContext())
306 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
307}
308
309static bool
310checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
311 Expr *ZExpr,
312 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
313 if (S.DiagnoseUnexpandedParameterPack(E: XExpr) ||
314 (YExpr && S.DiagnoseUnexpandedParameterPack(E: YExpr)) ||
315 (ZExpr && S.DiagnoseUnexpandedParameterPack(E: 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(AI: Attr, Expr: Exprs[i], Val&: NumWG, Idx: i,
329 /*StrictlyUnsigned=*/true))
330 return true;
331 if (NumWG == 0) {
332 S.Diag(Loc: Attr.getLoc(), DiagID: diag::err_attribute_argument_is_zero)
333 << &Attr << Exprs[i]->getSourceRange();
334 return true;
335 }
336 }
337 }
338
339 return false;
340}
341
342AMDGPUMaxNumWorkGroupsAttr *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(S&: SemaRef, XExpr, YExpr, ZExpr,
348 Attr: TmpAttr))
349 return nullptr;
350
351 return ::new (Context)
352 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
353}
354
355void 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(A: Attr);
361}
362
363void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
364 const ParsedAttr &AL) {
365 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
366 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(Arg: 2) : nullptr;
367 addAMDGPUMaxNumWorkGroupsAttr(D, CI: AL, XExpr: AL.getArgAsExpr(Arg: 0), YExpr, ZExpr);
368}
369
370} // namespace clang
371