30 unsigned OrderIndex, ScopeIndex;
32 const auto *FD =
SemaRef.getCurFunctionDecl(
true);
33 assert(FD &&
"AMDGPU builtins should not be used outside of a function");
34 llvm::StringMap<bool> CallerFeatureMap;
40 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_load_lds:
41 case AMDGPU::BI__builtin_amdgcn_struct_ptr_buffer_load_lds:
42 case AMDGPU::BI__builtin_amdgcn_load_to_lds:
43 case AMDGPU::BI__builtin_amdgcn_global_load_lds: {
44 constexpr const int SizeIdx = 2;
48 SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);
50 switch (Size.getSExtValue()) {
63 diag::err_amdgcn_load_lds_size_invalid_value)
66 diag::note_amdgcn_load_lds_size_valid_value)
71 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
72 case AMDGPU::BI__builtin_amdgcn_set_fpenv:
74 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
75 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
76 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
77 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
81 case AMDGPU::BI__builtin_amdgcn_fence:
85 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
87 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
89 case AMDGPU::BI__builtin_amdgcn_update_dpp:
91 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp8:
92 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp8:
93 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_bf8:
94 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_bf8:
95 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f16_fp4:
96 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_bf16_fp4:
97 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp8:
98 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_bf8:
99 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk8_f32_fp4:
100 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_fp6:
101 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_fp6:
102 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f16_bf6:
103 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_bf16_bf6:
104 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_fp6:
105 case AMDGPU::BI__builtin_amdgcn_cvt_scale_pk16_f32_bf6:
106 return SemaRef.BuiltinConstantArgRange(TheCall, 2, 0, 15);
107 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
108 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
109 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
111 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
112 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
113 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
115 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
116 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
117 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
118 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
119 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
120 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
121 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
122 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
123 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
124 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
125 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
126 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
127 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
128 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
129 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
130 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
131 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
132 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
133 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
134 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
135 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
136 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
137 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
138 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
139 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
140 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
141 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
142 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
143 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
144 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
145 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
146 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
147 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
148 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
149 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
150 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
151 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
152 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
153 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
154 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
155 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
156 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32: {
157 StringRef FeatureList(
162 << FD->getDeclName() << FeatureList;
166 unsigned ArgCount = TheCall->
getNumArgs() - 1;
171 (
SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1),
Result));
173 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
174 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
175 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
176 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
177 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
178 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
179 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
180 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
181 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
182 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
183 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
184 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
185 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
186 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
187 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
188 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
189 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
190 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
191 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
192 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
193 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
194 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
195 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
196 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
197 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
198 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
199 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
200 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
201 StringRef FeatureList(
206 << FD->getDeclName() << FeatureList;
210 unsigned ArgCount = TheCall->
getNumArgs() - 1;
215 (
SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1),
Result));
222 auto ArgExpr = Arg.
get();
226 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)
227 << ArgExpr->getType();
228 auto Ord = ArgResult.
Val.
getInt().getZExtValue();
232 if (!llvm::isValidAtomicOrderingCABI(Ord))
233 return Diag(ArgExpr->getBeginLoc(),
234 diag::warn_atomic_op_has_invalid_memory_order)
235 << 0 << ArgExpr->getSourceRange();
236 switch (
static_cast<llvm::AtomicOrderingCABI
>(Ord)) {
237 case llvm::AtomicOrderingCABI::relaxed:
238 case llvm::AtomicOrderingCABI::consume:
239 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
240 return Diag(ArgExpr->getBeginLoc(),
241 diag::warn_atomic_op_has_invalid_memory_order)
242 << 0 << ArgExpr->getSourceRange();
244 case llvm::AtomicOrderingCABI::acquire:
245 case llvm::AtomicOrderingCABI::release:
246 case llvm::AtomicOrderingCABI::acq_rel:
247 case llvm::AtomicOrderingCABI::seq_cst:
251 Arg = TheCall->
getArg(ScopeIndex);
255 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1,
getASTContext()))
256 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)
257 << ArgExpr->getType();