clang 22.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
18#include "clang/Sema/Sema.h"
19#include "llvm/Support/AMDGPUAddrSpace.h"
20#include "llvm/Support/AtomicOrdering.h"
21#include <cstdint>
22
23namespace clang {
24
26
28 CallExpr *TheCall) {
29 // position of memory order and scope arguments in the builtin
30 unsigned OrderIndex, ScopeIndex;
31
32 const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
33 assert(FD && "AMDGPU builtins should not be used outside of a function");
34 llvm::StringMap<bool> CallerFeatureMap;
35 getASTContext().getFunctionFeatureMap(CallerFeatureMap, FD);
36 bool HasGFX950Insts =
37 Builtin::evaluateRequiredTargetFeatures("gfx950-insts", CallerFeatureMap);
38
39 switch (BuiltinID) {
40 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_lds:
41 case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_lds:
42 case AMDGPU::BI__builtin_amdgcn_load_to_lds:
43 case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
44 constexpr const int SizeIdx = 2;
45 llvm::APSInt Size;
46 Expr *ArgExpr = TheCall->getArg(SizeIdx);
47 [[maybe_unused]] ExprResult R =
48 SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
49 assert(!R.isInvalid());
50 switch (Size.getSExtValue()) {
51 case 1:
52 case 2:
53 case 4:
54 return false;
55 case 12:
56 case 16: {
57 if (HasGFX950Insts)
58 return false;
59 [[fallthrough]];
60 }
61 default:
62 SemaRef.targetDiag(ArgExpr->getExprLoc(),
63 diag::err_amdgcn_load_lds_size_invalid_value)
64 << ArgExpr->getSourceRange();
65 SemaRef.targetDiag(ArgExpr->getExprLoc(),
66 diag::note_amdgcn_load_lds_size_valid_value)
67 << HasGFX950Insts << ArgExpr->getSourceRange();
68 return true;
69 }
70 }
71 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
72 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
73 return false;
74 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
75 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
76 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
77 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
78 OrderIndex = 2;
79 ScopeIndex = 3;
80 break;
81 case AMDGPU::BI__builtin_amdgcn_fence:
82 OrderIndex = 0;
83 ScopeIndex = 1;
84 break;
85 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
86 return checkMovDPPFunctionCall(TheCall, 5, 1);
87 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
88 return checkMovDPPFunctionCall(TheCall, 2, 1);
89 case AMDGPU::BI__builtin_amdgcn_update_dpp:
90 return checkMovDPPFunctionCall(TheCall, 6, 2);
91 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp8:
92 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp8:
93 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_bf8:
94 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_bf8:
95 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp4:
96 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp4:
97 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp8:
98 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_bf8:
99 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp4:
100 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_fp6:
101 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_fp6:
102 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_bf6:
103 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
104 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
105 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
106 return SemaRef.BuiltinConstantArgRange(TheCall, 2, 0, 15);
107 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
108 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
109 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
110 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/false);
111 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
112 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
113 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
114 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
115 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
116 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
117 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
118 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
119 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
120 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
121 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
122 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
123 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
124 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
125 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
126 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
127 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
128 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
129 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
130 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
131 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
132 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
133 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
134 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
135 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
136 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
137 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
138 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
139 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
140 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
141 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
142 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
143 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
144 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
145 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
146 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
147 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
148 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
149 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
150 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
151 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
152 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
153 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
154 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
155 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
156 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32: {
157 StringRef FeatureList(
158 getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
160 CallerFeatureMap)) {
161 Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
162 << FD->getDeclName() << FeatureList;
163 return false;
164 }
165
166 unsigned ArgCount = TheCall->getNumArgs() - 1;
167 llvm::APSInt Result;
168
169 return (SemaRef.BuiltinConstantArg(TheCall, 0, Result)) ||
170 (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
171 (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
172 }
173 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
174 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
175 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
176 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
177 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
178 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
179 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
180 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
181 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
182 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
183 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
184 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
185 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
186 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
187 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
188 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
189 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
190 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
191 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
192 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
193 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
194 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
195 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
196 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
197 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
198 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
199 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
200 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
201 StringRef FeatureList(
202 getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
204 CallerFeatureMap)) {
205 Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
206 << FD->getDeclName() << FeatureList;
207 return false;
208 }
209
210 unsigned ArgCount = TheCall->getNumArgs() - 1;
211 llvm::APSInt Result;
212
213 return (SemaRef.BuiltinConstantArg(TheCall, 1, Result)) ||
214 (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
215 (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
216 }
217 default:
218 return false;
219 }
220
221 ExprResult Arg = TheCall->getArg(OrderIndex);
222 auto ArgExpr = Arg.get();
223 Expr::EvalResult ArgResult;
224
225 if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))
226 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
227 << ArgExpr->getType();
228 auto Ord = ArgResult.Val.getInt().getZExtValue();
229
230 // Check validity of memory ordering as per C11 / C++11's memory model.
231 // Only fence needs check. Atomic dec/inc allow all memory orders.
232 if (!llvm::isValidAtomicOrderingCABI(Ord))
233 return Diag(ArgExpr->getBeginLoc(),
234 diag::warn_atomic_op_has_invalid_memory_order)
235 << 0 << ArgExpr->getSourceRange();
236 switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
237 case llvm::AtomicOrderingCABI::relaxed:
238 case llvm::AtomicOrderingCABI::consume:
239 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
240 return Diag(ArgExpr->getBeginLoc(),
241 diag::warn_atomic_op_has_invalid_memory_order)
242 << 0 << ArgExpr->getSourceRange();
243 break;
244 case llvm::AtomicOrderingCABI::acquire:
245 case llvm::AtomicOrderingCABI::release:
246 case llvm::AtomicOrderingCABI::acq_rel:
247 case llvm::AtomicOrderingCABI::seq_cst:
248 break;
249 }
250
251 Arg = TheCall->getArg(ScopeIndex);
252 ArgExpr = Arg.get();
253 Expr::EvalResult ArgResult1;
254 // Check that sync scope is a constant literal
255 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))
256 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
257 << ArgExpr->getType();
258
259 return false;
260}
261
263 bool Fail = false;
264
265 // First argument is a global or generic pointer.
266 Expr *PtrArg = TheCall->getArg(0);
267 QualType PtrTy = PtrArg->getType()->getPointeeType();
268 unsigned AS = getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace());
269 if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&
270 AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {
271 Fail = true;
272 Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as)
273 << PtrArg->getSourceRange();
274 }
275
276 // Check atomic ordering
277 Expr *AtomicOrdArg = TheCall->getArg(IsStore ? 2 : 1);
278 Expr::EvalResult AtomicOrdArgRes;
279 if (!AtomicOrdArg->EvaluateAsInt(AtomicOrdArgRes, getASTContext()))
280 llvm_unreachable("Intrinsic requires imm for atomic ordering argument!");
281 auto Ord =
282 llvm::AtomicOrderingCABI(AtomicOrdArgRes.Val.getInt().getZExtValue());
283
284 // Atomic ordering cannot be acq_rel in any case, acquire for stores or
285 // release for loads.
286 if (!llvm::isValidAtomicOrderingCABI((unsigned)Ord) ||
287 (Ord == llvm::AtomicOrderingCABI::acq_rel) ||
288 Ord == (IsStore ? llvm::AtomicOrderingCABI::acquire
289 : llvm::AtomicOrderingCABI::release)) {
290 return Diag(AtomicOrdArg->getBeginLoc(),
291 diag::warn_atomic_op_has_invalid_memory_order)
292 << 0 << AtomicOrdArg->getSourceRange();
293 }
294
295 // Last argument is a string literal
296 Expr *Arg = TheCall->getArg(TheCall->getNumArgs() - 1);
298 Fail = true;
299 Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)
300 << Arg->getSourceRange();
301 }
302
303 return Fail;
304}
305
306bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
307 unsigned NumDataArgs) {
308 assert(NumDataArgs <= 2);
309 if (SemaRef.checkArgCountRange(TheCall, NumArgs, NumArgs))
310 return true;
311 Expr *Args[2];
312 QualType ArgTys[2];
313 for (unsigned I = 0; I != NumDataArgs; ++I) {
314 Args[I] = TheCall->getArg(I);
315 ArgTys[I] = Args[I]->getType();
316 // TODO: Vectors can also be supported.
317 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
318 SemaRef.Diag(Args[I]->getBeginLoc(),
319 diag::err_typecheck_cond_expect_int_float)
320 << ArgTys[I] << Args[I]->getSourceRange();
321 return true;
322 }
323 }
324 if (NumDataArgs < 2)
325 return false;
326
327 if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1]))
328 return false;
329
330 if (((ArgTys[0]->isUnsignedIntegerType() &&
331 ArgTys[1]->isSignedIntegerType()) ||
332 (ArgTys[0]->isSignedIntegerType() &&
333 ArgTys[1]->isUnsignedIntegerType())) &&
334 getASTContext().getTypeSize(ArgTys[0]) ==
335 getASTContext().getTypeSize(ArgTys[1]))
336 return false;
337
338 SemaRef.Diag(Args[1]->getBeginLoc(),
339 diag::err_typecheck_call_different_arg_types)
340 << ArgTys[0] << ArgTys[1];
341 return true;
342}
343
344static bool
346 const AMDGPUFlatWorkGroupSizeAttr &Attr) {
347 // Accept template arguments for now as they depend on something else.
348 // We'll get to check them when they eventually get instantiated.
349 if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
350 return false;
351
352 uint32_t Min = 0;
353 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
354 return true;
355
356 uint32_t Max = 0;
357 if (!S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
358 return true;
359
360 if (Min == 0 && Max != 0) {
361 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
362 << &Attr << 0;
363 return true;
364 }
365 if (Min > Max) {
366 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
367 << &Attr << 1;
368 return true;
369 }
370
371 return false;
372}
373
374AMDGPUFlatWorkGroupSizeAttr *
376 Expr *MinExpr, Expr *MaxExpr) {
377 ASTContext &Context = getASTContext();
378 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
379
380 if (checkAMDGPUFlatWorkGroupSizeArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
381 return nullptr;
382 return ::new (Context)
383 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
384}
385
387 const AttributeCommonInfo &CI,
388 Expr *MinExpr, Expr *MaxExpr) {
389 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
390 D->addAttr(Attr);
391}
392
394 const ParsedAttr &AL) {
395 Expr *MinExpr = AL.getArgAsExpr(0);
396 Expr *MaxExpr = AL.getArgAsExpr(1);
397
398 addAMDGPUFlatWorkGroupSizeAttr(D, AL, MinExpr, MaxExpr);
399}
400
401static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
402 Expr *MaxExpr,
403 const AMDGPUWavesPerEUAttr &Attr) {
404 if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
405 (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
406 return true;
407
408 // Accept template arguments for now as they depend on something else.
409 // We'll get to check them when they eventually get instantiated.
410 if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
411 return false;
412
413 uint32_t Min = 0;
414 if (!S.checkUInt32Argument(Attr, MinExpr, Min, 0))
415 return true;
416
417 uint32_t Max = 0;
418 if (MaxExpr && !S.checkUInt32Argument(Attr, MaxExpr, Max, 1))
419 return true;
420
421 if (Min == 0 && Max != 0) {
422 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
423 << &Attr << 0;
424 return true;
425 }
426 if (Max != 0 && Min > Max) {
427 S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
428 << &Attr << 1;
429 return true;
430 }
431
432 return false;
433}
434
435AMDGPUWavesPerEUAttr *
437 Expr *MinExpr, Expr *MaxExpr) {
438 ASTContext &Context = getASTContext();
439 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
440
441 if (checkAMDGPUWavesPerEUArguments(SemaRef, MinExpr, MaxExpr, TmpAttr))
442 return nullptr;
443
444 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
445}
446
448 Expr *MinExpr, Expr *MaxExpr) {
449 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
450 D->addAttr(Attr);
451}
452
455 return;
456
457 Expr *MinExpr = AL.getArgAsExpr(0);
458 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
459
460 addAMDGPUWavesPerEUAttr(D, AL, MinExpr, MaxExpr);
461}
462
464 uint32_t NumSGPR = 0;
465 Expr *NumSGPRExpr = AL.getArgAsExpr(0);
466 if (!SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))
467 return;
468
469 D->addAttr(::new (getASTContext())
470 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
471}
472
474 uint32_t NumVGPR = 0;
475 Expr *NumVGPRExpr = AL.getArgAsExpr(0);
476 if (!SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))
477 return;
478
479 D->addAttr(::new (getASTContext())
480 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
481}
482
483static bool
485 Expr *ZExpr,
486 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
487 if (S.DiagnoseUnexpandedParameterPack(XExpr) ||
488 (YExpr && S.DiagnoseUnexpandedParameterPack(YExpr)) ||
489 (ZExpr && S.DiagnoseUnexpandedParameterPack(ZExpr)))
490 return true;
491
492 // Accept template arguments for now as they depend on something else.
493 // We'll get to check them when they eventually get instantiated.
494 if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
495 (ZExpr && ZExpr->isValueDependent()))
496 return false;
497
498 uint32_t NumWG = 0;
499 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
500 for (int i = 0; i < 3; i++) {
501 if (Exprs[i]) {
502 if (!S.checkUInt32Argument(Attr, Exprs[i], NumWG, i,
503 /*StrictlyUnsigned=*/true))
504 return true;
505 if (NumWG == 0) {
506 S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
507 << &Attr << Exprs[i]->getSourceRange();
508 return true;
509 }
510 }
511 }
512
513 return false;
514}
515
517 const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
518 ASTContext &Context = getASTContext();
519 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
520
521 if (checkAMDGPUMaxNumWorkGroupsArguments(SemaRef, XExpr, YExpr, ZExpr,
522 TmpAttr))
523 return nullptr;
524
525 return ::new (Context)
526 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
527}
528
530 const AttributeCommonInfo &CI,
531 Expr *XExpr, Expr *YExpr,
532 Expr *ZExpr) {
533 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
534 D->addAttr(Attr);
535}
536
538 const ParsedAttr &AL) {
539 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
540 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(2) : nullptr;
541 addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
542}
543
544} // namespace clang
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:489
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
Definition ASTContext.h:220
void getFunctionFeatureMap(llvm::StringMap< bool > &FeatureMap, const FunctionDecl *) const
unsigned getTargetAddressSpace(LangAS AS) const
PtrTy get() const
Definition Ownership.h:171
bool isInvalid() const
Definition Ownership.h:167
Attr - This represents one attribute.
Definition Attr.h:44
SourceLocation getLocation() const
Definition Attr.h:97
SourceLocation getLoc() const
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition Expr.h:2877
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition Expr.h:3081
SourceLocation getBeginLoc() const
Definition Expr.h:3211
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
Definition Expr.h:3068
Decl - This represents one declaration (or definition), e.g.
Definition DeclBase.h:86
void addAttr(Attr *A)
This represents one expression.
Definition Expr.h:112
bool EvaluateAsInt(EvalResult &Result, const ASTContext &Ctx, SideEffectsKind AllowSideEffects=SE_NoSideEffects, bool InConstantContext=false) const
EvaluateAsInt - Return true if this is a constant which we can fold and convert to an integer,...
bool isValueDependent() const
Determines whether the value of this expression depends on.
Definition Expr.h:177
Expr * IgnoreParenImpCasts() LLVM_READONLY
Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...
Definition Expr.cpp:3085
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
Definition Expr.cpp:273
QualType getType() const
Definition Expr.h:144
ParsedAttr - Represents a syntactic attribute.
Definition ParsedAttr.h:119
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
Definition ParsedAttr.h:371
Expr * getArgAsExpr(unsigned Arg) const
Definition ParsedAttr.h:383
bool checkAtLeastNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has at least as many args as Num.
bool checkAtMostNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has at most as many args as Num.
A (possibly-)qualified type.
Definition TypeBase.h:937
void handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, const ParsedAttr &AL)
void addAMDGPUFlatWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size attribute to a particular declar...
bool checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore)
void handleAMDGPUFlatWorkGroupSizeAttr(Decl *D, const ParsedAttr &AL)
void handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL)
AMDGPUMaxNumWorkGroupsAttr * CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr)
Create an AMDGPUMaxNumWorkGroupsAttr attribute.
AMDGPUWavesPerEUAttr * CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
Create an AMDGPUWavesPerEUAttr attribute.
void handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL)
AMDGPUFlatWorkGroupSizeAttr * CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
Create an AMDGPUWavesPerEUAttr attribute.
bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs, unsigned NumDataArgs)
void handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL)
bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)
void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max)
addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a particular declaration.
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...
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID, bool DeferHint=false)
Emit a diagnostic.
Definition SemaBase.cpp:61
SemaBase(Sema &S)
Definition SemaBase.cpp:7
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:854
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:4834
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
Definition Stmt.cpp:338
SourceLocation getBeginLoc() const LLVM_READONLY
Definition Stmt.cpp:350
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
Definition Type.cpp:752
bool evaluateRequiredTargetFeatures(llvm::StringRef RequiredFatures, const llvm::StringMap< bool > &TargetFetureMap)
Returns true if the required target features of a builtin function are enabled.
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
Definition Address.h:330
@ Result
The result type of a method or function.
Definition TypeBase.h:905
static bool checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr, Expr *ZExpr, const AMDGPUMaxNumWorkGroupsAttr &Attr)
static bool checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUFlatWorkGroupSizeAttr &Attr)
ActionResult< Expr * > ExprResult
Definition Ownership.h:249
static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, const AMDGPUWavesPerEUAttr &Attr)
EvalResult is a struct with detailed info about an evaluated expression.
Definition Expr.h:645
APValue Val
Val - This is the value the expression can be folded to.
Definition Expr.h:647