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