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/AST/Decl.h"
15#include "clang/AST/DynamicRecursiveASTVisitor.h"
16#include "clang/AST/Expr.h"
17#include "clang/Basic/DiagnosticFrontend.h"
18#include "clang/Basic/DiagnosticSema.h"
19#include "clang/Basic/TargetBuiltins.h"
20#include "clang/Basic/TargetInfo.h"
21#include "clang/Sema/Ownership.h"
22#include "clang/Sema/Scope.h"
23#include "clang/Sema/Sema.h"
24#include "llvm/ADT/SmallVector.h"
25#include "llvm/ADT/StringExtras.h"
26#include "llvm/ADT/StringMap.h"
27#include "llvm/Support/AMDGPUAddrSpace.h"
28#include "llvm/Support/AtomicOrdering.h"
29#include "llvm/TargetParser/AMDGPUTargetParser.h"
30#include "llvm/TargetParser/AtomicScope.h"
31#include <cstdint>
32#include <utility>
33
34namespace clang {
35
36SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {}
37
38bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI,
39 unsigned BuiltinID,
40 CallExpr *TheCall) {
41 const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
42 assert(FD && "AMDGPU builtins should not be used outside of a function");
43 llvm::StringMap<bool> CallerFeatureMap;
44 getASTContext().getFunctionFeatureMap(FeatureMap&: CallerFeatureMap, FD);
45 bool HasGFX950Insts =
46 Builtin::evaluateRequiredTargetFeatures(RequiredFatures: "gfx950-insts", TargetFetureMap: CallerFeatureMap);
47
48 switch (BuiltinID) {
49 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_lds:
50 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_async_lds:
51 case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_lds:
52 case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_async_lds:
53 case AMDGPU::BI__builtin_amdgcn_load_to_lds:
54 case AMDGPU::BI__builtin_amdgcn_load_async_to_lds:
55 case AMDGPU::BI__builtin_amdgcn_global_load_lds:
56 case AMDGPU::BI__builtin_amdgcn_global_load_async_lds: {
57 constexpr const int SizeIdx = 2;
58 llvm::APSInt Size;
59 Expr *ArgExpr = TheCall->getArg(Arg: SizeIdx);
60 // Check for instantiation-dependent expressions (e.g., involving template
61 // parameters). These will be checked again during template instantiation.
62 if (ArgExpr->isInstantiationDependent())
63 return false;
64 [[maybe_unused]] ExprResult R =
65 SemaRef.VerifyIntegerConstantExpression(E: ArgExpr, Result: &Size);
66 assert(!R.isInvalid());
67 switch (Size.getSExtValue()) {
68 case 1:
69 case 2:
70 case 4:
71 return false;
72 case 12:
73 case 16: {
74 if (HasGFX950Insts)
75 return false;
76 [[fallthrough]];
77 }
78 default:
79 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
80 DiagID: diag::err_amdgcn_load_lds_size_invalid_value)
81 << ArgExpr->getSourceRange();
82 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
83 DiagID: diag::note_amdgcn_load_lds_size_valid_value)
84 << HasGFX950Insts << ArgExpr->getSourceRange();
85 return true;
86 }
87 }
88 case AMDGPU::BI__builtin_amdgcn_uicmp:
89 case AMDGPU::BI__builtin_amdgcn_uicmpl:
90 case AMDGPU::BI__builtin_amdgcn_sicmp:
91 case AMDGPU::BI__builtin_amdgcn_sicmpl:
92 case AMDGPU::BI__builtin_amdgcn_fcmp:
93 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
94 // These builtins are deprecated in favor of
95 // __builtin_amdgcn_ballot_{w32|w64}. Suggest the replacement matching the
96 // wavefront size of the calling function.
97 bool IsWave32 = Builtin::evaluateRequiredTargetFeatures(RequiredFatures: "wavefrontsize32",
98 TargetFetureMap: CallerFeatureMap);
99 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::warn_deprecated_builtin)
100 << getASTContext().BuiltinInfo.getQuotedName(ID: BuiltinID)
101 << (IsWave32 ? "__builtin_amdgcn_ballot_w32"
102 : "__builtin_amdgcn_ballot_w64");
103 return false;
104 }
105 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
106 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
107 return false;
108 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
109 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
110 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
111 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
112 case AMDGPU::BI__builtin_amdgcn_fence: {
113 bool IsFence = BuiltinID == AMDGPU::BI__builtin_amdgcn_fence;
114 unsigned OrderIndex = IsFence ? 0 : 2;
115 unsigned ScopeIndex = IsFence ? 1 : 3;
116 Expr *OrderExpr = TheCall->getArg(Arg: OrderIndex);
117 Expr *ScopeExpr = TheCall->getArg(Arg: ScopeIndex);
118
119 // Checks requiring constant evaluation are deferred until instantiation.
120 if (OrderExpr->isInstantiationDependent() ||
121 ScopeExpr->isInstantiationDependent())
122 return false;
123
124 Expr::EvalResult OrderResult;
125 if (!OrderExpr->EvaluateAsInt(Result&: OrderResult, Ctx: getASTContext()))
126 return Diag(Loc: OrderExpr->getExprLoc(), DiagID: diag::err_typecheck_expect_int)
127 << OrderExpr->getType();
128 uint64_t Ord = OrderResult.Val.getInt().getZExtValue();
129
130 // Check validity of memory ordering as per C11 / C++11's memory model.
131 // Only fence needs check. Atomic dec/inc allow all memory orders.
132 if (!llvm::isValidAtomicOrderingCABI(I: Ord))
133 return Diag(Loc: OrderExpr->getBeginLoc(),
134 DiagID: diag::warn_atomic_op_has_invalid_memory_order)
135 << 0 << OrderExpr->getSourceRange();
136 switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
137 case llvm::AtomicOrderingCABI::relaxed:
138 case llvm::AtomicOrderingCABI::consume:
139 if (IsFence)
140 return Diag(Loc: OrderExpr->getBeginLoc(),
141 DiagID: diag::warn_atomic_op_has_invalid_memory_order)
142 << 0 << OrderExpr->getSourceRange();
143 break;
144 case llvm::AtomicOrderingCABI::acquire:
145 case llvm::AtomicOrderingCABI::release:
146 case llvm::AtomicOrderingCABI::acq_rel:
147 case llvm::AtomicOrderingCABI::seq_cst:
148 break;
149 }
150
151 Expr::EvalResult ScopeResult;
152 // Check that sync scope is a constant literal
153 if (!ScopeExpr->EvaluateAsConstantExpr(Result&: ScopeResult, Ctx: getASTContext()))
154 return Diag(Loc: ScopeExpr->getExprLoc(), DiagID: diag::err_expr_not_string_literal)
155 << ScopeExpr->getType();
156
157 // Reject any string that is not a valid synchronization scope name.
158 std::optional<std::string> ScopeName =
159 ScopeExpr->tryEvaluateString(Ctx&: getASTContext());
160 const llvm::Triple &TT = TI.getTriple();
161 if (ScopeName && !llvm::parseAtomicScopeIRString(T: TT, Name: *ScopeName)) {
162 return Diag(Loc: ScopeExpr->getExprLoc(), DiagID: diag::err_invalid_sync_scope)
163 << *ScopeName << ScopeExpr->getSourceRange();
164 }
165
166 return false;
167 }
168 case AMDGPU::BI__builtin_amdgcn_s_setreg:
169 return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
170 /*High=*/UINT16_MAX);
171 case AMDGPU::BI__builtin_amdgcn_s_wait_event: {
172 llvm::APSInt Result;
173 if (SemaRef.BuiltinConstantArg(TheCall, ArgNum: 0, Result))
174 return true;
175
176 bool IsGFX12Plus = Builtin::evaluateRequiredTargetFeatures(
177 RequiredFatures: "gfx12-insts", TargetFetureMap: CallerFeatureMap);
178
179 // gfx11 -> gfx12 changed the interpretation of the bitmask. gfx12 inverted
180 // the intepretation for export_ready, but shifted the used bit by 1. Thus
181 // waiting for the export_ready event can use a value of 2 universally.
182 if (((IsGFX12Plus && !Result[1]) || (!IsGFX12Plus && Result[0])) ||
183 Result.getZExtValue() > 2) {
184 Expr *ArgExpr = TheCall->getArg(Arg: 0);
185 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
186 DiagID: diag::warn_amdgpu_s_wait_event_mask_no_effect_target)
187 << ArgExpr->getSourceRange();
188 SemaRef.targetDiag(Loc: ArgExpr->getExprLoc(),
189 DiagID: diag::note_amdgpu_s_wait_event_suggested_value)
190 << ArgExpr->getSourceRange();
191 }
192
193 return false;
194 }
195 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
196 return checkMovDPPFunctionCall(TheCall, NumArgs: 5, NumDataArgs: 1);
197 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
198 return checkMovDPPFunctionCall(TheCall, NumArgs: 2, NumDataArgs: 1);
199 case AMDGPU::BI__builtin_amdgcn_update_dpp:
200 return checkMovDPPFunctionCall(TheCall, NumArgs: 6, NumDataArgs: 2);
201 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp8:
202 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp8:
203 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_bf8:
204 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_bf8:
205 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp4:
206 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp4:
207 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp8:
208 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_bf8:
209 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp4:
210 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_fp6:
211 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_fp6:
212 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_bf6:
213 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
214 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
215 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
216 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_bf16_bf6:
217 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_bf16_fp6:
218 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f16_bf6:
219 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f16_fp6:
220 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f32_bf6:
221 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk32_f32_fp6:
222 return SemaRef.BuiltinConstantArgRange(TheCall, ArgNum: 2, Low: 0, High: 15);
223 case AMDGPU::BI__builtin_amdgcn_av_load_b128:
224 return checkAVLoadStore(TheCall, /*IsStore=*/false);
225 case AMDGPU::BI__builtin_amdgcn_av_store_b128:
226 return checkAVLoadStore(TheCall, /*IsStore=*/true);
227 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
228 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
229 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
230 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/false);
231 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
232 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
233 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
234 return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
235 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
236 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
237 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128:
238 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
239 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
240 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
241 return checkAtomicMonitorLoad(TheCall);
242 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
243 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
244 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
245 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
246 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
247 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
248 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
249 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
250 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
251 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
252 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
253 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
254 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
255 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
256 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
257 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
258 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
259 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
260 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
261 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
262 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
263 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
264 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
265 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
266 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
267 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
268 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
269 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
270 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
271 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
272 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
273 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
274 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
275 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
276 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
277 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
278 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
279 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
280 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
281 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
282 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
283 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
284 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
285 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
286 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
287 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
288 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
289 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
290 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
291 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
292 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
293 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
294 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
295 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
296 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
297 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
298 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
299 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
300 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
301 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
302 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
303 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
304 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
305 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
306 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
307 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
308 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
309 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
310 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
311 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
312 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
313 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
314 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
315 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
316 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
317 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
318 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
319 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
320 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
321 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
322 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
323 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
324 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
325 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32: {
326 StringRef FeatureList(
327 getASTContext().BuiltinInfo.getRequiredFeatures(ID: BuiltinID));
328 if (!Builtin::evaluateRequiredTargetFeatures(RequiredFatures: FeatureList,
329 TargetFetureMap: CallerFeatureMap)) {
330 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_builtin_needs_feature)
331 << FD->getDeclName() << FeatureList;
332 return false;
333 }
334
335 unsigned ArgCount = TheCall->getNumArgs() - 1;
336 llvm::APSInt Result;
337
338 // Compilain about dmask values which are too huge to fully fit into 4 bits
339 // (which is the actual size of the dmask in corresponding HW instructions).
340 constexpr unsigned DMaskArgNo = 0;
341 constexpr int Low = 0;
342 constexpr int High = 15;
343 if (SemaRef.BuiltinConstantArg(TheCall, ArgNum: DMaskArgNo, Result) ||
344 SemaRef.BuiltinConstantArgRange(TheCall, ArgNum: DMaskArgNo, Low, High,
345 /* RangeIsError = */ true))
346 return true;
347
348 // Dmask indicates which elements should be returned and it is not possible
349 // to return more values than there are elements in return type.
350 int NumElementsInRetTy = 1;
351 const Type *RetTy = TheCall->getType().getTypePtr();
352 if (auto *VTy = dyn_cast<VectorType>(Val: RetTy))
353 NumElementsInRetTy = VTy->getNumElements();
354 int NumActiveBitsInDMask =
355 llvm::popcount(Value: static_cast<uint8_t>(Result.getExtValue()));
356 if (NumActiveBitsInDMask > NumElementsInRetTy) {
357 Diag(Loc: TheCall->getBeginLoc(),
358 DiagID: diag::err_amdgcn_dmask_has_too_many_bits_set);
359 return true;
360 }
361
362 // For gather, only one bit can be set indicating which exact component to
363 // return.
364 bool ExtraGatherChecks =
365 (BuiltinID ==
366 AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32 ||
367 BuiltinID ==
368 AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32) &&
369 SemaRef.BuiltinConstantArgPower2(TheCall, ArgNum: 0);
370
371 return ExtraGatherChecks ||
372 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: ArgCount, Result)) ||
373 (SemaRef.BuiltinConstantArg(TheCall, ArgNum: (ArgCount - 1), Result));
374 }
375 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
376 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
377 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
378 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
379 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
380 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
381 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
382 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
383 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
384 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
385 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
386 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
387 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
388 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
389 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
390 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
391 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
392 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
393 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
394 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
395 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
396 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
397 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
398 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
399 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
400 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
401 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
402 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
403 StringRef FeatureList(
404 getASTContext().BuiltinInfo.getRequiredFeatures(ID: BuiltinID));
405 if (!Builtin::evaluateRequiredTargetFeatures(RequiredFatures: FeatureList,
406 TargetFetureMap: CallerFeatureMap)) {
407 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_builtin_needs_feature)
408 << FD->getDeclName() << FeatureList;
409 return false;
410 }
411
412 unsigned ArgCount = TheCall->getNumArgs() - 1;
413 llvm::APSInt Result;
414
415 // Complain about dmask values which are too huge to fully fit into 4 bits
416 // (which is the actual size of the dmask in corresponding HW instructions).
417 constexpr unsigned DMaskArgNo = 1;
418 return SemaRef.BuiltinConstantArgRange(TheCall, ArgNum: DMaskArgNo, /*Low=*/0,
419 /*High=*/15,
420 /*RangeIsError=*/true) ||
421 SemaRef.BuiltinConstantArg(TheCall, ArgNum: ArgCount, Result) ||
422 SemaRef.BuiltinConstantArg(TheCall, ArgNum: (ArgCount - 1), Result);
423 }
424 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
425 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
426 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8) {
427 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: 7, MaxArgCount: 8))
428 return true;
429 if (TheCall->getNumArgs() == 7)
430 return false;
431 } else if (BuiltinID ==
432 AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8) {
433 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: 8, MaxArgCount: 9))
434 return true;
435 if (TheCall->getNumArgs() == 8)
436 return false;
437 }
438 // Check if the last argument (clamp operand) is a constant and is
439 // convertible to bool.
440 Expr *ClampArg = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
441 // 1) Ensure clamp argument is a constant expression
442 llvm::APSInt ClampValue;
443 if (!SemaRef.VerifyIntegerConstantExpression(E: ClampArg, Result: &ClampValue)
444 .isUsable())
445 return true;
446 // 2) Check if the argument can be converted to bool type
447 if (!SemaRef.Context.hasSameType(T1: ClampArg->getType(),
448 T2: SemaRef.Context.BoolTy)) {
449 // Try to convert to bool
450 QualType BoolTy = SemaRef.Context.BoolTy;
451 ExprResult ClampExpr(ClampArg);
452 SemaRef.CheckSingleAssignmentConstraints(LHSType: BoolTy, RHS&: ClampExpr);
453 if (ClampExpr.isInvalid())
454 return true;
455 }
456 return false;
457 }
458 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
459 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
460 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
461 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
462 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
463 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
464 return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
465 /*High=*/0) ||
466 SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/2, /*Low=*/0,
467 /*High=*/0);
468 default:
469 return false;
470 }
471}
472
473bool SemaAMDGPU::checkAtomicOrderingCABIArg(Expr *E, bool MayLoad,
474 bool MayStore) {
475 Expr::EvalResult AtomicOrdArgRes;
476 if (!E->EvaluateAsInt(Result&: AtomicOrdArgRes, Ctx: getASTContext()))
477 llvm_unreachable("Intrinsic requires imm for atomic ordering argument!");
478 auto Ord =
479 llvm::AtomicOrderingCABI(AtomicOrdArgRes.Val.getInt().getZExtValue());
480
481 // Atomic ordering cannot be acq_rel in any case, acquire for stores or
482 // release for loads.
483 if (!llvm::isValidAtomicOrderingCABI(I: (unsigned)Ord) ||
484 (!(MayLoad && MayStore) && (Ord == llvm::AtomicOrderingCABI::acq_rel)) ||
485 (!MayLoad && Ord == llvm::AtomicOrderingCABI::acquire) ||
486 (!MayStore && Ord == llvm::AtomicOrderingCABI::release)) {
487 return Diag(Loc: E->getBeginLoc(), DiagID: diag::warn_atomic_op_has_invalid_memory_order)
488 << 0 << E->getSourceRange();
489 }
490
491 return false;
492}
493
494// Check that the first argument to TheCall is a global or generic pointer.
495static bool checkGlobalOrFlatPointerArg(SemaAMDGPU &S, CallExpr *TheCall) {
496 Expr *PtrArg = TheCall->getArg(Arg: 0);
497 QualType PtrTy = PtrArg->getType()->getPointeeType();
498 unsigned AS =
499 S.getASTContext().getTargetAddressSpace(AS: PtrTy.getAddressSpace());
500 if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&
501 AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {
502 return S.Diag(Loc: TheCall->getBeginLoc(),
503 DiagID: diag::err_amdgcn_global_or_flat_pointer_required)
504 << PtrArg->getSourceRange();
505 }
506 return false;
507}
508
509static bool checkScopeAsInt(SemaAMDGPU &S, Expr *Scope) {
510 if (Scope->isValueDependent())
511 return false;
512 auto ScopeModel = AtomicScopeModel::create(K: AtomicScopeModelKind::Generic);
513 if (std::optional<llvm::APSInt> Result =
514 Scope->getIntegerConstantExpr(Ctx: S.SemaRef.Context)) {
515 if (!ScopeModel->isValid(S: Result->getZExtValue())) {
516 return S.Diag(Loc: Scope->getBeginLoc(),
517 DiagID: diag::err_atomic_op_has_invalid_sync_scope)
518 << Scope->getSourceRange();
519 }
520 }
521 return false;
522}
523
524bool SemaAMDGPU::checkAVLoadStore(CallExpr *TheCall, bool IsStore) {
525 if (checkGlobalOrFlatPointerArg(S&: *this, TheCall))
526 return true;
527
528 Expr *Scope = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
529 return checkScopeAsInt(S&: *this, Scope);
530}
531
532bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) {
533 bool Fail = checkGlobalOrFlatPointerArg(S&: *this, TheCall);
534
535 Expr *AO = TheCall->getArg(Arg: IsStore ? 2 : 1);
536 Expr *Scope = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
537
538 if (AO->isValueDependent() || Scope->isValueDependent())
539 return false;
540
541 // Check atomic ordering
542 Fail |=
543 checkAtomicOrderingCABIArg(E: TheCall->getArg(Arg: IsStore ? 2 : 1),
544 /*MayLoad=*/!IsStore, /*MayStore=*/IsStore);
545
546 // Last argument is the syncscope as a string literal.
547 if (!isa<StringLiteral>(Val: Scope->IgnoreParenImpCasts())) {
548 Diag(Loc: TheCall->getBeginLoc(), DiagID: diag::err_expr_not_string_literal)
549 << Scope->getSourceRange();
550 Fail = true;
551 }
552
553 return Fail;
554}
555
556bool SemaAMDGPU::checkAtomicMonitorLoad(CallExpr *TheCall) {
557 Expr *AO = TheCall->getArg(Arg: 1);
558 Expr *Scope = TheCall->getArg(Arg: TheCall->getNumArgs() - 1);
559
560 if (AO->isValueDependent() || Scope->isValueDependent())
561 return false;
562
563 bool Fail = checkAtomicOrderingCABIArg(E: AO, /*MayLoad=*/true,
564 /*MayStore=*/false);
565 Fail |= checkScopeAsInt(S&: *this, Scope);
566 return Fail;
567}
568
569bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
570 unsigned NumDataArgs) {
571 assert(NumDataArgs <= 2);
572 if (SemaRef.checkArgCountRange(Call: TheCall, MinArgCount: NumArgs, MaxArgCount: NumArgs))
573 return true;
574 Expr *Args[2];
575 QualType ArgTys[2];
576 for (unsigned I = 0; I != NumDataArgs; ++I) {
577 Args[I] = TheCall->getArg(Arg: I);
578 ArgTys[I] = Args[I]->getType();
579 // TODO: Vectors can also be supported.
580 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
581 SemaRef.Diag(Loc: Args[I]->getBeginLoc(),
582 DiagID: diag::err_typecheck_cond_expect_int_float)
583 << ArgTys[I] << Args[I]->getSourceRange();
584 return true;
585 }
586 }
587 if (NumDataArgs < 2)
588 return false;
589
590 if (getASTContext().hasSameUnqualifiedType(T1: ArgTys[0], T2: ArgTys[1]))
591 return false;
592
593 if (((ArgTys[0]->isUnsignedIntegerType() &&
594 ArgTys[1]->isSignedIntegerType()) ||
595 (ArgTys[0]->isSignedIntegerType() &&
596 ArgTys[1]->isUnsignedIntegerType())) &&
597 getASTContext().getTypeSize(T: ArgTys[0]) ==
598 getASTContext().getTypeSize(T: ArgTys[1]))
599 return false;
600
601 SemaRef.Diag(Loc: Args[1]->getBeginLoc(),
602 DiagID: diag::err_typecheck_call_different_arg_types)
603 << ArgTys[0] << ArgTys[1];
604 return true;
605}
606
607static bool
608checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
609 const AMDGPUFlatWorkGroupSizeAttr &Attr) {
610 // Accept template arguments for now as they depend on something else.
611 // We'll get to check them when they eventually get instantiated.
612 if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
613 return false;
614
615 uint32_t Min = 0;
616 if (!S.checkUInt32Argument(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
617 return true;
618
619 uint32_t Max = 0;
620 if (!S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
621 return true;
622
623 if (Min == 0 && Max != 0) {
624 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
625 << &Attr << 0;
626 return true;
627 }
628 if (Min > Max) {
629 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
630 << &Attr << 1;
631 return true;
632 }
633
634 return false;
635}
636
637AMDGPUFlatWorkGroupSizeAttr *
638SemaAMDGPU::CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI,
639 Expr *MinExpr, Expr *MaxExpr) {
640 ASTContext &Context = getASTContext();
641 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
642
643 if (checkAMDGPUFlatWorkGroupSizeArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
644 return nullptr;
645 return ::new (Context)
646 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);
647}
648
649void SemaAMDGPU::addAMDGPUFlatWorkGroupSizeAttr(Decl *D,
650 const AttributeCommonInfo &CI,
651 Expr *MinExpr, Expr *MaxExpr) {
652 if (auto *Attr = CreateAMDGPUFlatWorkGroupSizeAttr(CI, MinExpr, MaxExpr))
653 D->addAttr(A: Attr);
654}
655
656void SemaAMDGPU::handleAMDGPUFlatWorkGroupSizeAttr(Decl *D,
657 const ParsedAttr &AL) {
658 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
659 Expr *MaxExpr = AL.getArgAsExpr(Arg: 1);
660
661 addAMDGPUFlatWorkGroupSizeAttr(D, CI: AL, MinExpr, MaxExpr);
662}
663
664static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
665 Expr *MaxExpr,
666 const AMDGPUWavesPerEUAttr &Attr) {
667 if (S.DiagnoseUnexpandedParameterPack(E: MinExpr) ||
668 (MaxExpr && S.DiagnoseUnexpandedParameterPack(E: MaxExpr)))
669 return true;
670
671 // Accept template arguments for now as they depend on something else.
672 // We'll get to check them when they eventually get instantiated.
673 if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
674 return false;
675
676 uint32_t Min = 0;
677 if (!S.checkUInt32Argument(AI: Attr, Expr: MinExpr, Val&: Min, Idx: 0))
678 return true;
679
680 uint32_t Max = 0;
681 if (MaxExpr && !S.checkUInt32Argument(AI: Attr, Expr: MaxExpr, Val&: Max, Idx: 1))
682 return true;
683
684 if (Min == 0 && Max != 0) {
685 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
686 << &Attr << 0;
687 return true;
688 }
689 if (Max != 0 && Min > Max) {
690 S.Diag(Loc: Attr.getLocation(), DiagID: diag::err_attribute_argument_invalid)
691 << &Attr << 1;
692 return true;
693 }
694
695 return false;
696}
697
698AMDGPUWavesPerEUAttr *
699SemaAMDGPU::CreateAMDGPUWavesPerEUAttr(const AttributeCommonInfo &CI,
700 Expr *MinExpr, Expr *MaxExpr) {
701 ASTContext &Context = getASTContext();
702 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);
703
704 if (checkAMDGPUWavesPerEUArguments(S&: SemaRef, MinExpr, MaxExpr, Attr: TmpAttr))
705 return nullptr;
706
707 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);
708}
709
710void SemaAMDGPU::addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI,
711 Expr *MinExpr, Expr *MaxExpr) {
712 if (auto *Attr = CreateAMDGPUWavesPerEUAttr(CI, MinExpr, MaxExpr))
713 D->addAttr(A: Attr);
714}
715
716void SemaAMDGPU::handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL) {
717 if (!AL.checkAtLeastNumArgs(S&: SemaRef, Num: 1) || !AL.checkAtMostNumArgs(S&: SemaRef, Num: 2))
718 return;
719
720 Expr *MinExpr = AL.getArgAsExpr(Arg: 0);
721 Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
722
723 addAMDGPUWavesPerEUAttr(D, CI: AL, MinExpr, MaxExpr);
724}
725
726void SemaAMDGPU::handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL) {
727 Diag(Loc: AL.getLoc(), DiagID: diag::warn_amdgpu_num_reg_attr_deprecated) << AL;
728
729 uint32_t NumSGPR = 0;
730 Expr *NumSGPRExpr = AL.getArgAsExpr(Arg: 0);
731 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumSGPRExpr, Val&: NumSGPR))
732 return;
733
734 D->addAttr(A: ::new (getASTContext())
735 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));
736}
737
738void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
739 Diag(Loc: AL.getLoc(), DiagID: diag::warn_amdgpu_num_reg_attr_deprecated) << AL;
740
741 uint32_t NumVGPR = 0;
742 Expr *NumVGPRExpr = AL.getArgAsExpr(Arg: 0);
743 if (!SemaRef.checkUInt32Argument(AI: AL, Expr: NumVGPRExpr, Val&: NumVGPR))
744 return;
745
746 D->addAttr(A: ::new (getASTContext())
747 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
748}
749
750static bool
751checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
752 Expr *ZExpr,
753 const AMDGPUMaxNumWorkGroupsAttr &Attr) {
754 if (S.DiagnoseUnexpandedParameterPack(E: XExpr) ||
755 (YExpr && S.DiagnoseUnexpandedParameterPack(E: YExpr)) ||
756 (ZExpr && S.DiagnoseUnexpandedParameterPack(E: ZExpr)))
757 return true;
758
759 // Accept template arguments for now as they depend on something else.
760 // We'll get to check them when they eventually get instantiated.
761 if (XExpr->isValueDependent() || (YExpr && YExpr->isValueDependent()) ||
762 (ZExpr && ZExpr->isValueDependent()))
763 return false;
764
765 uint32_t NumWG = 0;
766 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};
767 for (int i = 0; i < 3; i++) {
768 if (Exprs[i]) {
769 if (!S.checkUInt32Argument(AI: Attr, Expr: Exprs[i], Val&: NumWG, Idx: i,
770 /*StrictlyUnsigned=*/true))
771 return true;
772 if (NumWG == 0) {
773 S.Diag(Loc: Attr.getLoc(), DiagID: diag::err_attribute_argument_is_zero)
774 << &Attr << Exprs[i]->getSourceRange();
775 return true;
776 }
777 }
778 }
779
780 return false;
781}
782
783AMDGPUMaxNumWorkGroupsAttr *SemaAMDGPU::CreateAMDGPUMaxNumWorkGroupsAttr(
784 const AttributeCommonInfo &CI, Expr *XExpr, Expr *YExpr, Expr *ZExpr) {
785 ASTContext &Context = getASTContext();
786 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);
787 assert(!SemaRef.isSFINAEContext() &&
788 "Can't produce SFINAE diagnostic pointing to temporary attribute");
789
790 if (checkAMDGPUMaxNumWorkGroupsArguments(S&: SemaRef, XExpr, YExpr, ZExpr,
791 Attr: TmpAttr))
792 return nullptr;
793
794 return ::new (Context)
795 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);
796}
797
798void SemaAMDGPU::addAMDGPUMaxNumWorkGroupsAttr(Decl *D,
799 const AttributeCommonInfo &CI,
800 Expr *XExpr, Expr *YExpr,
801 Expr *ZExpr) {
802 if (auto *Attr = CreateAMDGPUMaxNumWorkGroupsAttr(CI, XExpr, YExpr, ZExpr))
803 D->addAttr(A: Attr);
804}
805
806void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
807 const ParsedAttr &AL) {
808 Expr *YExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(Arg: 1) : nullptr;
809 Expr *ZExpr = (AL.getNumArgs() > 2) ? AL.getArgAsExpr(Arg: 2) : nullptr;
810 addAMDGPUMaxNumWorkGroupsAttr(D, CI: AL, XExpr: AL.getArgAsExpr(Arg: 0), YExpr, ZExpr);
811}
812
813Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) {
814 CallExpr *CE = cast<CallExpr>(Val: E->IgnoreParens());
815 ASTContext &Ctx = getASTContext();
816 QualType BoolTy = Ctx.getLogicalOperationType();
817 SourceLocation Loc = CE->getExprLoc();
818
819 if (!CE->getBuiltinCallee())
820 return *ExpandedPredicates
821 .insert(Ptr: SemaRef.BuildBoolLiteral(Loc, Value: false).get())
822 .first;
823
824 bool P = false;
825 unsigned BI = CE->getBuiltinCallee();
826 if (Ctx.BuiltinInfo.isAuxBuiltinID(ID: BI))
827 BI = Ctx.BuiltinInfo.getAuxBuiltinID(ID: BI);
828
829 if (BI == AMDGPU::BI__builtin_amdgcn_processor_is) {
830 auto *GFX = dyn_cast<StringLiteral>(Val: CE->getArg(Arg: 0)->IgnoreParenCasts());
831 if (!GFX) {
832 Diag(Loc, DiagID: diag::err_amdgcn_processor_is_arg_not_literal);
833 return nullptr;
834 }
835
836 StringRef N = GFX->getString();
837 const TargetInfo &TI = Ctx.getTargetInfo();
838 if (llvm::AMDGPU::parseArchAMDGCN(CPU: N) == llvm::AMDGPU::GK_NONE) {
839 Diag(Loc, DiagID: diag::err_amdgcn_processor_is_arg_invalid_value) << N;
840 SmallVector<StringRef, 64> ValidList;
841 llvm::AMDGPU::fillValidArchListAMDGCN(Values&: ValidList);
842 if (!ValidList.empty())
843 Diag(Loc, DiagID: diag::note_amdgcn_processor_is_valid_options)
844 << llvm::join(R&: ValidList, Separator: ", ");
845 return nullptr;
846 }
847 if (TI.getTriple().isSPIRV()) {
848 CE->setType(BoolTy);
849 return *ExpandedPredicates.insert(Ptr: CE).first;
850 }
851
852 P = TI.isProcessorName(Name: N);
853 } else {
854 Expr *Arg = CE->getArg(Arg: 0);
855 if (!Arg || Arg->getType() != Ctx.BuiltinFnTy) {
856 Diag(Loc, DiagID: diag::err_amdgcn_is_invocable_arg_invalid_value) << Arg;
857 return nullptr;
858 }
859
860 if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
861 CE->setType(BoolTy);
862 return *ExpandedPredicates.insert(Ptr: CE).first;
863 }
864
865 auto *FD = cast<FunctionDecl>(Val: Arg->getReferencedDeclOfCallee());
866
867 StringRef RF = Ctx.BuiltinInfo.getRequiredFeatures(ID: FD->getBuiltinID());
868 llvm::StringMap<bool> CF;
869 Ctx.getFunctionFeatureMap(FeatureMap&: CF, FD);
870
871 P = Builtin::evaluateRequiredTargetFeatures(RequiredFatures: RF, TargetFetureMap: CF);
872 }
873
874 return *ExpandedPredicates.insert(Ptr: SemaRef.BuildBoolLiteral(Loc, Value: P).get())
875 .first;
876}
877
878bool SemaAMDGPU::IsPredicate(Expr *E) const {
879 return ExpandedPredicates.contains(Ptr: E);
880}
881
882void SemaAMDGPU::AddPotentiallyUnguardedBuiltinUser(FunctionDecl *FD) {
883 PotentiallyUnguardedBuiltinUsers.insert(Ptr: FD);
884}
885
886bool SemaAMDGPU::HasPotentiallyUnguardedBuiltinUsage(FunctionDecl *FD) const {
887 return PotentiallyUnguardedBuiltinUsers.contains(Ptr: FD);
888}
889
890namespace {
891/// This class implements -Wamdgpu-unguarded-builtin-usage.
892///
893/// This is done with a traversal of the AST of a function that includes a
894/// call to a target specific builtin. Whenever we encounter an \c if of the
895/// form: \c if(__builtin_amdgcn_is_invocable), we consider the then statement
896/// guarded.
897class DiagnoseUnguardedBuiltins : public DynamicRecursiveASTVisitor {
898 // TODO: this could eventually be extended to consider what happens when there
899 // are multiple target architectures specified via target("arch=gfxXXX")
900 // target("arch=gfxyyy") etc., as well as feature disabling via "-XXX".
901 Sema &SemaRef;
902
903 SmallVector<StringRef> TargetFeatures;
904 SmallVector<std::pair<SourceLocation, StringRef>> CurrentGFXIP;
905 SmallVector<unsigned> GuardedBuiltins;
906
907 static Expr *FindPredicate(Expr *Cond) {
908 if (auto *CE = dyn_cast<CallExpr>(Val: Cond)) {
909 if (CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_is_invocable ||
910 CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is)
911 return Cond;
912 } else if (auto *UO = dyn_cast<UnaryOperator>(Val: Cond)) {
913 return FindPredicate(Cond: UO->getSubExpr());
914 } else if (auto *BO = dyn_cast<BinaryOperator>(Val: Cond)) {
915 if ((Cond = FindPredicate(Cond: BO->getLHS())))
916 return Cond;
917 return FindPredicate(Cond: BO->getRHS());
918 }
919 return nullptr;
920 }
921
922 bool EnterPredicateGuardedContext(CallExpr *P);
923 void ExitPredicateGuardedContext(bool WasProcessorCheck);
924 bool TraverseGuardedStmt(Stmt *S, CallExpr *P);
925
926public:
927 DiagnoseUnguardedBuiltins(Sema &SemaRef) : SemaRef(SemaRef) {
928 if (auto *TAT = SemaRef.getCurFunctionDecl(AllowLambda: true)->getAttr<TargetAttr>()) {
929 // We use the somewhat misnamed x86 accessors because they provide exactly
930 // what we require.
931 TAT->getX86AddedFeatures(Out&: TargetFeatures);
932 if (auto GFXIP = TAT->getX86Architecture())
933 CurrentGFXIP.emplace_back(Args: TAT->getLocation(), Args&: *GFXIP);
934 }
935 }
936
937 bool TraverseLambdaExpr(LambdaExpr *LE) override {
938 if (SemaRef.AMDGPU().HasPotentiallyUnguardedBuiltinUsage(
939 FD: LE->getCallOperator()))
940 return true; // We have already handled this.
941 return DynamicRecursiveASTVisitor::TraverseLambdaExpr(S: LE);
942 }
943
944 bool TraverseStmt(Stmt *S) override {
945 if (!S)
946 return true;
947 return DynamicRecursiveASTVisitor::TraverseStmt(S);
948 }
949
950 void IssueDiagnostics(Stmt *S) { TraverseStmt(S); }
951
952 bool TraverseIfStmt(IfStmt *If) override {
953 if (auto *CE = dyn_cast_or_null<CallExpr>(Val: FindPredicate(Cond: If->getCond())))
954 return TraverseGuardedStmt(S: If, P: CE);
955 return DynamicRecursiveASTVisitor::TraverseIfStmt(S: If);
956 }
957
958 bool TraverseCaseStmt(CaseStmt *CS) override {
959 return TraverseStmt(S: CS->getSubStmt());
960 }
961
962 bool TraverseConditionalOperator(ConditionalOperator *CO) override {
963 if (auto *CE = dyn_cast_or_null<CallExpr>(Val: FindPredicate(Cond: CO->getCond())))
964 return TraverseGuardedStmt(S: CO, P: CE);
965 return DynamicRecursiveASTVisitor::TraverseConditionalOperator(S: CO);
966 }
967
968 bool VisitAsmStmt(AsmStmt *ASM) override;
969 bool VisitCallExpr(CallExpr *CE) override;
970};
971
972bool DiagnoseUnguardedBuiltins::EnterPredicateGuardedContext(CallExpr *P) {
973 bool IsProcessorCheck =
974 P->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is;
975
976 if (IsProcessorCheck) {
977 StringRef G = cast<clang::StringLiteral>(Val: P->getArg(Arg: 0))->getString();
978 // TODO: handle generic ISAs.
979 if (!CurrentGFXIP.empty() && G != CurrentGFXIP.back().second) {
980 SemaRef.Diag(Loc: P->getExprLoc(),
981 DiagID: diag::err_amdgcn_conflicting_is_processor_options)
982 << P;
983 SemaRef.Diag(Loc: CurrentGFXIP.back().first,
984 DiagID: diag::note_amdgcn_previous_is_processor_guard);
985 }
986 CurrentGFXIP.emplace_back(Args: P->getExprLoc(), Args&: G);
987 } else {
988 auto *FD = cast<FunctionDecl>(
989 Val: cast<DeclRefExpr>(Val: P->getArg(Arg: 0))->getReferencedDeclOfCallee());
990 GuardedBuiltins.push_back(Elt: FD->getBuiltinID());
991 }
992
993 return IsProcessorCheck;
994}
995
996void DiagnoseUnguardedBuiltins::ExitPredicateGuardedContext(bool WasProcCheck) {
997 if (WasProcCheck)
998 CurrentGFXIP.pop_back();
999 else
1000 GuardedBuiltins.pop_back();
1001}
1002
1003inline std::pair<Stmt *, Stmt *> GetTraversalOrder(Stmt *S) {
1004 std::pair<Stmt *, Stmt *> Ordered;
1005 Expr *Condition = nullptr;
1006
1007 if (auto *CO = dyn_cast<ConditionalOperator>(Val: S)) {
1008 Condition = CO->getCond();
1009 Ordered = {CO->getTrueExpr(), CO->getFalseExpr()};
1010 } else if (auto *If = dyn_cast<IfStmt>(Val: S)) {
1011 Condition = If->getCond();
1012 Ordered = {If->getThen(), If->getElse()};
1013 }
1014
1015 if (auto *UO = dyn_cast<UnaryOperator>(Val: Condition))
1016 if (UO->getOpcode() == UnaryOperatorKind::UO_LNot)
1017 std::swap(a&: Ordered.first, b&: Ordered.second);
1018
1019 return Ordered;
1020}
1021
1022bool DiagnoseUnguardedBuiltins::TraverseGuardedStmt(Stmt *S, CallExpr *P) {
1023 assert(S && "Unexpected missing Statement!");
1024 assert(P && "Unexpected missing Predicate!");
1025
1026 auto [Guarded, Unguarded] = GetTraversalOrder(S);
1027
1028 bool WasProcessorCheck = EnterPredicateGuardedContext(P);
1029
1030 bool Continue = TraverseStmt(S: Guarded);
1031
1032 ExitPredicateGuardedContext(WasProcCheck: WasProcessorCheck);
1033
1034 return Continue && TraverseStmt(S: Unguarded);
1035}
1036
1037bool DiagnoseUnguardedBuiltins::VisitAsmStmt(AsmStmt *ASM) {
1038 // TODO: should we check if the ASM is valid for the target? Can we?
1039 if (!CurrentGFXIP.empty())
1040 return true;
1041
1042 std::string S = ASM->generateAsmString(C: SemaRef.getASTContext());
1043 SemaRef.Diag(Loc: ASM->getAsmLoc(), DiagID: diag::warn_amdgcn_unguarded_asm_stmt) << S;
1044 SemaRef.Diag(Loc: ASM->getAsmLoc(), DiagID: diag::note_amdgcn_unguarded_asm_silence) << S;
1045
1046 return true;
1047}
1048
1049bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) {
1050 unsigned ID = CE->getBuiltinCallee();
1051 Builtin::Context &BInfo = SemaRef.getASTContext().BuiltinInfo;
1052
1053 if (!ID)
1054 return true;
1055 if (!BInfo.isTSBuiltin(ID))
1056 return true;
1057 if (ID == AMDGPU::BI__builtin_amdgcn_processor_is ||
1058 ID == AMDGPU::BI__builtin_amdgcn_is_invocable)
1059 return true;
1060 if (llvm::find(Range&: GuardedBuiltins, Val: ID) != GuardedBuiltins.end())
1061 return true;
1062
1063 StringRef FL(BInfo.getRequiredFeatures(ID));
1064 llvm::StringMap<bool> FeatureMap;
1065 if (CurrentGFXIP.empty()) {
1066 for (auto &&F : TargetFeatures)
1067 FeatureMap[F] = true;
1068 for (auto &&GID : GuardedBuiltins)
1069 for (auto &&F : llvm::split(Str: BInfo.getRequiredFeatures(ID: GID), Separator: ','))
1070 FeatureMap[F] = true;
1071 } else {
1072 static const llvm::Triple AMDGCN(llvm::Triple::amdgpu,
1073 llvm::Triple::NoSubArch, llvm::Triple::AMD,
1074 llvm::Triple::AMDHSA);
1075 llvm::AMDGPU::fillAMDGPUFeatureMap(GPU: CurrentGFXIP.back().second, T: AMDGCN,
1076 Features&: FeatureMap);
1077 }
1078
1079 FunctionDecl *BI = CE->getDirectCallee();
1080 SourceLocation BICallLoc = CE->getExprLoc();
1081 if (Builtin::evaluateRequiredTargetFeatures(RequiredFatures: FL, TargetFetureMap: FeatureMap)) {
1082 SemaRef.Diag(Loc: BICallLoc, DiagID: diag::warn_amdgcn_unguarded_builtin) << BI;
1083 SemaRef.Diag(Loc: BICallLoc, DiagID: diag::note_amdgcn_unguarded_builtin_silence) << BI;
1084 } else {
1085 StringRef GFXIP = CurrentGFXIP.empty() ? "" : CurrentGFXIP.back().second;
1086 SemaRef.Diag(Loc: BICallLoc, DiagID: diag::err_amdgcn_incompatible_builtin)
1087 << BI << FL << !CurrentGFXIP.empty() << GFXIP;
1088 if (!CurrentGFXIP.empty())
1089 SemaRef.Diag(Loc: CurrentGFXIP.back().first,
1090 DiagID: diag::note_amdgcn_previous_is_processor_guard);
1091 }
1092
1093 return true;
1094}
1095} // Unnamed namespace
1096
1097void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) {
1098 DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(S: FD->getBody());
1099}
1100
1101bool SemaAMDGPU::checkAMDGPUTypeSupport(QualType Ty, SourceLocation Loc) {
1102 ASTContext &Ctx = getASTContext();
1103 llvm::Triple TT = Ctx.getTargetInfo().getTriple();
1104 const Type *BaseTy = Ty->getPointeeOrArrayElementType();
1105
1106 if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
1107 // The AMDGPU named barrier type requires special handling in the back-end
1108 // and is not supported for SPIR-V
1109 if (BaseTy->isAMDGPUNamedBarrierType()) {
1110 SemaRef.Diag(Loc, DiagID: diag::err_amdgpu_target_ext_type_unsupported)
1111 << Ty << TT.str();
1112 return false;
1113 }
1114 }
1115
1116 return true;
1117}
1118
1119static FieldDecl *getNamedBarrierField(const RecordDecl *R) {
1120 for (FieldDecl *FD : R->fields()) {
1121 QualType FDTy = FD->getType();
1122 if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper())
1123 return FD;
1124 }
1125
1126 return nullptr;
1127}
1128
1129void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) {
1130 ASTContext &Context = getASTContext();
1131 if (R->isInvalidDecl())
1132 return;
1133
1134 if (!Context.getTargetInfo().hasAMDGPUTypes() &&
1135 (!Context.getAuxTargetInfo() ||
1136 !Context.getAuxTargetInfo()->hasAMDGPUTypes()))
1137 return;
1138
1139 bool IsWrapper = false;
1140 std::function<void()> DiagWrapperNote;
1141
1142 // First, check if this is a named barrier wrapper by virtue of the class
1143 // declaring a named barrier field. This covers both C and C++.
1144 if (FieldDecl *NamedBarrField = getNamedBarrierField(R)) {
1145 // If this record contains a named barrier field, it must have only one
1146 // field.
1147 if (R->getNumFields() > 1) {
1148 SemaRef.Diag(Loc: NamedBarrField->getLocation(),
1149 DiagID: diag::err_amdgpu_invalid_field_not_a_wrapper)
1150 << NamedBarrField->getType();
1151 SemaRef.Diag(
1152 Loc: R->getLocation(),
1153 DiagID: diag::note_amdgpu_not_a_named_barrier_wrapper_too_many_fields)
1154 << R->getDeclName();
1155 return;
1156 }
1157
1158 IsWrapper = true;
1159 DiagWrapperNote = [this, R, NamedBarrField]() {
1160 SemaRef.Diag(Loc: NamedBarrField->getLocation(),
1161 DiagID: diag::note_amdgpu_named_barrier_reason_field)
1162 << R->getDeclName() << NamedBarrField->getDeclName();
1163 };
1164 }
1165
1166 // Then, for C++ classes, check if this is a named barrier wrapper by virtue
1167 // of inheriting one.
1168 const auto *CxxR = dyn_cast<CXXRecordDecl>(Val: R);
1169 if (CxxR && !IsWrapper) {
1170 for (CXXBaseSpecifier BS : CxxR->bases()) {
1171 const RecordDecl *Base = BS.getType()->getAsRecordDecl();
1172 if (!Base || !Base->hasAttr<AMDGPUNamedBarrierWrapperAttr>())
1173 continue;
1174
1175 IsWrapper = true;
1176 DiagWrapperNote = [this, BS, R]() {
1177 // Print using the CXXBaseSpecifier type as it includes the template
1178 // parameters.
1179 SemaRef.Diag(Loc: BS.getBeginLoc(),
1180 DiagID: diag::note_amdgpu_named_barrier_reason_inherited)
1181 << R->getDeclName() << BS.getType();
1182 };
1183 }
1184 }
1185
1186 if (!IsWrapper)
1187 return;
1188
1189 // Set the attribute even if the wrapper may be found to be invalid later.
1190 R->addAttr(
1191 A: AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Ctx&: Context, Range: SourceRange()));
1192
1193 // This is a wrapper CXXRecordDecl, it must have a C++11 standard layout.
1194 if (CxxR && !CxxR->isCXX11StandardLayout()) {
1195 SemaRef.Diag(Loc: R->getLocation(),
1196 DiagID: diag::err_amdgpu_named_barrier_wrapper_non_standard_layout)
1197 << R->getDeclName();
1198 assert(DiagWrapperNote &&
1199 "IsWrapper is set but no context diagnostic provided");
1200 DiagWrapperNote();
1201 }
1202}
1203} // namespace clang
1204