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/DiagnosticFrontend.h"
15#include "clang/Basic/DiagnosticSema.h"
16#include "clang/Basic/TargetBuiltins.h"
17#include "clang/Sema/Ownership.h"
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
25SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
26
27bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
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(FeatureMap&: CallerFeatureMap, FD);
36 bool HasGFX950Insts =
37 Builtin::evaluateRequiredTargetFeatures(RequiredFatures: "gfx950-insts", TargetFetureMap: 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(Arg: SizeIdx);
47 // Check for instantiation-dependent expressions (e.g., involving template
48 // parameters). These will be checked again during template instantiation.
49 if (ArgExpr->isInstantiationDependent())
50 return false;
51 [[maybe_unused]] ExprResult R =
52 SemaRef.VerifyIntegerConstantExpression(E: ArgExpr, Result: &Size);
53 assert(!R.isInvalid());
54 switch (Size.getSExtValue()) {
55 case 1:
56 case 2:
57 case 4:
58 return false;
59 case 12:
60 case 16: {
61 if (HasGFX950Insts)
62 return false;
63 [[fallthrough]];
64 }
65 default:
66 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
67 DiagID: diag::err_amdgcn_load_lds_size_invalid_value)
68 << ArgExpr->getSourceRange();
69 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
70 DiagID: diag::note_amdgcn_load_lds_size_valid_value)
71 << HasGFX950Insts << ArgExpr->getSourceRange();
72 return true;
73 }
74 }
75 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
76 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
77 return false;
78 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
79 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
80 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
81 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
82 OrderIndex = 2;
83 ScopeIndex = 3;
84 break;
85 case AMDGPU::BI__builtin_amdgcn_fence:
86 OrderIndex = 0;
87 ScopeIndex = 1;
88 break;
89 case AMDGPU::BI__builtin_amdgcn_s_setreg:
90 return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
91 /*High=*/UINT16_MAX);
92 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
93 return checkMovDPPFunctionCall(TheCall, NumArgs: 5, NumDataArgs: 1);
94 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
95 return checkMovDPPFunctionCall(TheCall, NumArgs: 2, NumDataArgs: 1);
96 case AMDGPU::BI__builtin_amdgcn_update_dpp:
97 return checkMovDPPFunctionCall(TheCall, NumArgs: 6, NumDataArgs: 2);
98 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp8:
99 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp8:
100 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_bf8:
101 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_bf8:
102 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp4:
103 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp4:
104 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp8:
105 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_bf8:
106 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp4:
107 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_fp6:
108 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_fp6:
109 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_bf6:
110 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
111 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
112 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
113 return SemaRef.BuiltinConstantArgRange(TheCall, ArgNum: 2, Low: 0, High: 15);
114 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
115 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
116 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
117 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/false);
118 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
119 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
120 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
121 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
122 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
123 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
124 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
125 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
126 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
127 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
128 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
129 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
130 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
131 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
132 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
133 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
134 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
135 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
136 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
137 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
138 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
139 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
140 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
141 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
142 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
143 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
144 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
145 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
146 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
147 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
148 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
149 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
150 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
151 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
152 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
153 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
154 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
155 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
156 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
157 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
158 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
159 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
160 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
161 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
162 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
163 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
164 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
165 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
166 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
167 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
168 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
169 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
170 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
171 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
172 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
173 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
174 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
175 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
176 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
177 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
178 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
179 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
180 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
181 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
182 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
183 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
184 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
185 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
186 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
187 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
188 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
189 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
190 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
191 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
192 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
193 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
194 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
195 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
196 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
197 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
198 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
199 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
200 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
201 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
202 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
203 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
204 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {
205 StringRef FeatureList(
206 getASTContext().BuiltinInfo.getRequiredFeatures(ID: BuiltinID));
207 if (!Builtin::evaluateRequiredTargetFeatures(RequiredFatures: FeatureList,
208 TargetFetureMap: CallerFeatureMap)) {
209 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_builtin_needs_feature)
210 << FD->getDeclName() << FeatureList;
211 return false;
212 }
213
214 unsigned ArgCount = TheCall->getNumArgs() - 1;
215 llvm::APSInt Result;
216
217 return (SemaRef.BuiltinConstantArg(TheCall, ArgNum: 0, Result)) ||
218 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: ArgCount, Result)) ||
219 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: (ArgCount - 1), Result));
220 }
221 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
222 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
223 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
224 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
225 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
226 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
227 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
228 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
229 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
230 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
231 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
232 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
233 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
234 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
235 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
236 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
237 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
238 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
239 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
240 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
241 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
242 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
243 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
244 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
245 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
246 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
247 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
248 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
249 StringRef FeatureList(
250 getASTContext().BuiltinInfo.getRequiredFeatures(ID: BuiltinID));
251 if (!Builtin::evaluateRequiredTargetFeatures(RequiredFatures: FeatureList,
252 TargetFetureMap: CallerFeatureMap)) {
253 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_builtin_needs_feature)
254 << FD->getDeclName() << FeatureList;
255 return false;
256 }
257
258 unsigned ArgCount = TheCall->getNumArgs() - 1;
259 llvm::APSInt Result;
260
261 return (SemaRef.BuiltinConstantArg(TheCall, ArgNum: 1, Result)) ||
262 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: ArgCount, Result)) ||
263 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: (ArgCount - 1), Result));
264 }
265 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
266 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
267 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8) {
268 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: 7, MaxArgCount: 8))
269 return true;
270 if (TheCall->getNumArgs() == 7)
271 return false;
272 } else if (BuiltinID ==
273 AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8) {
274 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: 8, MaxArgCount: 9))
275 return true;
276 if (TheCall->getNumArgs() == 8)
277 return false;
278 }
279 // Check if the last argument (clamp operand) is a constant and is
280 // convertible to bool.
281 Expr *ClampArg = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
282 // 1) Ensure clamp argument is a constant expression
283 llvm::APSInt ClampValue;
284 if (!SemaRef.VerifyIntegerConstantExpression(E: ClampArg, Result: &ClampValue)
285 .isUsable())
286 return true;
287 // 2) Check if the argument can be converted to bool type
288 if (!SemaRef.Context.hasSameType(T1: ClampArg->getType(),
289 T2: SemaRef.Context.BoolTy)) {
290 // Try to convert to bool
291 QualType BoolTy = SemaRef.Context.BoolTy;
292 ExprResult ClampExpr(ClampArg);
293 SemaRef.CheckSingleAssignmentConstraints(LHSType: BoolTy, RHS&: ClampExpr);
294 if (ClampExpr.isInvalid())
295 return true;
296 }
297 return false;
298 }
299 default:
300 return false;
301 }
302
303 ExprResult Arg = TheCall->getArg(Arg: OrderIndex);
304 auto ArgExpr = Arg.get();
305 Expr::EvalResult ArgResult;
306
307 if (!ArgExpr->EvaluateAsInt(Result&: ArgResult, Ctx: getASTContext()))
308 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: diag::err_typecheck_expect_int)
309 << ArgExpr->getType();
310 auto Ord = ArgResult.Val.getInt().getZExtValue();
311
312 // Check validity of memory ordering as per C11 / C++11's memory model.
313 // Only fence needs check. Atomic dec/inc allow all memory orders.
314 if (!llvm::isValidAtomicOrderingCABI(I: Ord))
315 return Diag(Loc: ArgExpr->getBeginLoc(),
316 DiagID: diag::warn_atomic_op_has_invalid_memory_order)
317 << 0 << ArgExpr->getSourceRange();
318 switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
319 case llvm::AtomicOrderingCABI::relaxed:
320 case llvm::AtomicOrderingCABI::consume:
321 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
322 return Diag(Loc: ArgExpr->getBeginLoc(),
323 DiagID: diag::warn_atomic_op_has_invalid_memory_order)
324 << 0 << ArgExpr->getSourceRange();
325 break;
326 case llvm::AtomicOrderingCABI::acquire:
327 case llvm::AtomicOrderingCABI::release:
328 case llvm::AtomicOrderingCABI::acq_rel:
329 case llvm::AtomicOrderingCABI::seq_cst:
330 break;
331 }
332
333 Arg = TheCall->getArg(Arg: ScopeIndex);
334 ArgExpr = Arg.get();
335 Expr::EvalResult ArgResult1;
336 // Check that sync scope is a constant literal
337 if (!ArgExpr->EvaluateAsConstantExpr(Result&: ArgResult1, Ctx: getASTContext()))
338 return Diag(Loc: ArgExpr->getExprLoc(), DiagID: diag::err_expr_not_string_literal)
339 << ArgExpr->getType();
340
341 return false;
342}
343
344bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) {
345 bool Fail = false;
346
347 // First argument is a global or generic pointer.
348 Expr *PtrArg = TheCall->getArg(Arg: 0);
349 QualType PtrTy = PtrArg->getType()->getPointeeType();
350 unsigned AS = getASTContext().getTargetAddressSpace(AS: PtrTy.getAddressSpace());
351 if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&
352 AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {
353 Fail = true;
354 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_amdgcn_coop_atomic_invalid_as)
355 << PtrArg->getSourceRange();
356 }
357
358 // Check atomic ordering
359 Expr *AtomicOrdArg = TheCall->getArg(Arg: IsStore ? 2 : 1);
360 Expr::EvalResult AtomicOrdArgRes;
361 if (!AtomicOrdArg->EvaluateAsInt(Result&: AtomicOrdArgRes, Ctx: getASTContext()))
362 llvm_unreachable("Intrinsic requires imm for atomic ordering argument!");
363 auto Ord =
364 llvm::AtomicOrderingCABI(AtomicOrdArgRes.Val.getInt().getZExtValue());
365
366 // Atomic ordering cannot be acq_rel in any case, acquire for stores or
367 // release for loads.
368 if (!llvm::isValidAtomicOrderingCABI(I: (unsigned)Ord) ||
369 (Ord == llvm::AtomicOrderingCABI::acq_rel) ||
370 Ord == (IsStore ? llvm::AtomicOrderingCABI::acquire
371 : llvm::AtomicOrderingCABI::release)) {
372 return Diag(Loc: AtomicOrdArg->getBeginLoc(),
373 DiagID: diag::warn_atomic_op_has_invalid_memory_order)
374 << 0 << AtomicOrdArg->getSourceRange();
375 }
376
377 // Last argument is a string literal
378 Expr *Arg = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
379 if (!isa<StringLiteral>(Val: Arg->IgnoreParenImpCasts())) {
380 Fail = true;
381 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_expr_not_string_literal)
382 << Arg->getSourceRange();
383 }
384
385 return Fail;
386}
387
388bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
389 unsigned NumDataArgs) {
390 assert(NumDataArgs <= 2);
391 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: NumArgs, MaxArgCount: NumArgs))
392 return true;
393 Expr *Args[2];
394 QualType ArgTys[2];
395 for (unsigned I = 0; I != NumDataArgs; ++I) {
396 Args[I] = TheCall->getArg(Arg: I);
397 ArgTys[I] = Args[I]->getType();
398 // TODO: Vectors can also be supported.
399 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
400 SemaRef.Diag(Loc: Args[I]->getBeginLoc(),
401 DiagID: diag::err_typecheck_cond_expect_int_float)
402 << ArgTys[I] << Args[I]->getSourceRange();
403 return true;
404 }
405 }
406 if (NumDataArgs < 2)
407 return false;
408
409 if (getASTContext().hasSameUnqualifiedType(T1: ArgTys[0], T2: ArgTys[1]))
410 return false;
411
412 if (((ArgTys[0]->isUnsignedIntegerType() &&
413 ArgTys[1]->isSignedIntegerType()) ||
414 (ArgTys[0]->isSignedIntegerType() &&
415 ArgTys[1]->isUnsignedIntegerType())) &&
416 getASTContext().getTypeSize(T: ArgTys[0]) ==
417 getASTContext().getTypeSize(T: ArgTys[1]))
418 return false;
419
420 SemaRef.Diag(Loc: Args[1]->getBeginLoc(),
421 DiagID: diag::err_typecheck_call_different_arg_types)
422 << ArgTys[0] << ArgTys[1];
423 return true;
424}
425
426static bool
427checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
428 const AMDGPUFlatWorkGroupSizeAttr &Attr) {
429 // Accept template arguments for now as they depend on something else.
430 // We'll get to check them when they eventually get instantiated.
431 if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
432 return false;
433
434 uint32_t Min = 0;
435 if (!S.checkUInt32Argument(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
436 return true;
437
438 uint32_t Max = 0;
439 if (!S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
440 return true;
441
442 if (Min == 0 && Max != 0) {
443 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
444 << &Attr << 0;
445 return true;
446 }
447 if (Min > Max) {
448 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
449 << &Attr << 1;
450 return true;
451 }
452
453 return false;
454}
455
456AMDGPUFlatWorkGroupSizeAttr *
457SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
458 Expr *MinExpr, Expr *MaxExpr) {
459 ASTContext &Context = getASTContext();
460 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
461
462 if (checkAMDGPUFlatWorkGroupSizeArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
463 return nullptr;
464 return ::new (Context)
465 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
466}
467
468void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
469 const AttributeCommonInfo &CI,
470 Expr *MinExpr, Expr *MaxExpr) {
471 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
472 D->addAttr(A: Attr);
473}
474
475void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
476 const ParsedAttr &AL) {
477 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
478 Expr *MaxExpr = AL.getArgAsExpr(Arg: 1);
479
480 addAMDGPUFlatWorkGroupSizeAttr(D, CI: AL, MinExpr, MaxExpr);
481}
482
483static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
484 Expr *MaxExpr,
485 const AMDGPUWavesPerEUAttr &Attr) {
486 if (S.DiagnoseUnexpandedParameterPack(E: MinExpr) ||
487 (MaxExpr && S.DiagnoseUnexpandedParameterPack(E: MaxExpr)))
488 return true;
489
490 // Accept template arguments for now as they depend on something else.
491 // We'll get to check them when they eventually get instantiated.
492 if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
493 return false;
494
495 uint32_t Min = 0;
496 if (!S.checkUInt32Argument(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
497 return true;
498
499 uint32_t Max = 0;
500 if (MaxExpr && !S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
501 return true;
502
503 if (Min == 0 && Max != 0) {
504 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
505 << &Attr << 0;
506 return true;
507 }
508 if (Max != 0 && Min > Max) {
509 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
510 << &Attr << 1;
511 return true;
512 }
513
514 return false;
515}
516
517AMDGPUWavesPerEUAttr *
518SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
519 Expr *MinExpr, Expr *MaxExpr) {
520 ASTContext &Context = getASTContext();
521 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
522
523 if (checkAMDGPUWavesPerEUArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
524 return nullptr;
525
526 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
527}
528
529void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
530 Expr *MinExpr, Expr *MaxExpr) {
531 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
532 D->addAttr(A: Attr);
533}
534
535void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
536 if (!AL.checkAtLeastNumArgs(S&: SemaRef, Num: 1) || !AL.checkAtMostNumArgs(S&: SemaRef, Num: 2))
537 return;
538
539 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
540 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
541
542 addAMDGPUWavesPerEUAttr(D, CI: AL, MinExpr, MaxExpr);
543}
544
545void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
546 uint32_t NumSGPR = 0;
547 Expr *NumSGPRExpr = AL.getArgAsExpr(Arg: 0);
548 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumSGPRExpr, Val&: NumSGPR))
549 return;
550
551 D->addAttr(A: ::new (getASTContext())
552 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
553}
554
555void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
556 uint32_t NumVGPR = 0;
557 Expr *NumVGPRExpr = AL.getArgAsExpr(Arg: 0);
558 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumVGPRExpr, Val&: NumVGPR))
559 return;
560
561 D->addAttr(A: ::new (getASTContext())
562 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
563}
564
565static bool
566checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
567 Expr *ZExpr,
568 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
569 if (S.DiagnoseUnexpandedParameterPack(E: XExpr) ||
570 (YExpr && S.DiagnoseUnexpandedParameterPack(E: YExpr)) ||
571 (ZExpr && S.DiagnoseUnexpandedParameterPack(E: ZExpr)))
572 return true;
573
574 // Accept template arguments for now as they depend on something else.
575 // We'll get to check them when they eventually get instantiated.
576 if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
577 (ZExpr && ZExpr->isValueDependent()))
578 return false;
579
580 uint32_t NumWG = 0;
581 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
582 for (int i = 0; i < 3; i++) {
583 if (Exprs[i]) {
584 if (!S.checkUInt32Argument(AI: Attr, Expr: Exprs[i], Val&: NumWG, Idx: i,
585 /*StrictlyUnsigned=*/true))
586 return true;
587 if (NumWG == 0) {
588 S.Diag(Loc: Attr.getLoc(), DiagID: diag::err_attribute_argument_is_zero)
589 << &Attr << Exprs[i]->getSourceRange();
590 return true;
591 }
592 }
593 }
594
595 return false;
596}
597
598AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
599 const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
600 ASTContext &Context = getASTContext();
601 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
602 assert(!SemaRef.isSFINAEContext() &&
603 "Can't produce SFINAE diagnostic pointing to temporary attribute");
604
605 if (checkAMDGPUMaxNumWorkGroupsArguments(S&: SemaRef, XExpr, YExpr, ZExpr,
606 Attr: TmpAttr))
607 return nullptr;
608
609 return ::new (Context)
610 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
611}
612
613void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
614 const AttributeCommonInfo &CI,
615 Expr *XExpr, Expr *YExpr,
616 Expr *ZExpr) {
617 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
618 D->addAttr(A: Attr);
619}
620
621void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
622 const ParsedAttr &AL) {
623 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
624 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(Arg: 2) : nullptr;
625 addAMDGPUMaxNumWorkGroupsAttr(D, CI: AL, XExpr: AL.getArgAsExpr(Arg: 0), YExpr, ZExpr);
626}
627
628} // namespace clang
629