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