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