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