clang: lib/Sema/SemaAMDGPU.cpp Source File (original) (raw)

29

30 unsigned OrderIndex, ScopeIndex;

31

32 const auto *FD = SemaRef.getCurFunctionDecl(true);

33 assert(FD && "AMDGPU builtins should not be used outside of a function");

34 llvm::StringMap CallerFeatureMap;

36 bool HasGFX950Insts =

38

39 switch (BuiltinID) {

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;

45 llvm::APSInt Size;

46 Expr *ArgExpr = TheCall->getArg(SizeIdx);

48 SemaRef.VerifyIntegerConstantExpression(ArgExpr, &Size);

50 switch (Size.getSExtValue()) {

51 case 1:

52 case 2:

53 case 4:

54 return false;

55 case 12:

56 case 16: {

57 if (HasGFX950Insts)

58 return false;

59 [[fallthrough]];

60 }

61 default:

63 diag::err_amdgcn_load_lds_size_invalid_value)

66 diag::note_amdgcn_load_lds_size_valid_value)

68 return true;

69 }

70 }

71 case AMDGPU::BI__builtin_amdgcn_get_fpenv:

72 case AMDGPU::BI__builtin_amdgcn_set_fpenv:

73 return false;

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:

78 OrderIndex = 2;

79 ScopeIndex = 3;

80 break;

81 case AMDGPU::BI__builtin_amdgcn_fence:

82 OrderIndex = 0;

83 ScopeIndex = 1;

84 break;

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 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:

158 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:

159 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:

160 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:

161 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:

162 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:

163 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:

164 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:

165 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:

166 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:

167 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:

168 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:

169 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:

170 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:

171 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:

172 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:

173 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:

174 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:

175 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:

176 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:

177 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:

178 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:

179 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:

180 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:

181 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:

182 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:

183 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:

184 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:

185 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:

186 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:

187 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:

188 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:

189 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:

190 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:

191 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:

192 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:

193 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:

194 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:

195 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:

196 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:

197 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {

198 StringRef FeatureList(

199 getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));

201 CallerFeatureMap)) {

202 Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)

203 << FD->getDeclName() << FeatureList;

204 return false;

205 }

206

207 unsigned ArgCount = TheCall->getNumArgs() - 1;

209

210 return (SemaRef.BuiltinConstantArg(TheCall, 0, Result)) ||

211 (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||

212 (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));

213 }

214 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:

215 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:

216 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:

217 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:

218 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:

219 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:

220 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:

221 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:

222 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:

223 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:

224 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:

225 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:

226 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:

227 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:

228 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:

229 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:

230 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:

231 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:

232 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:

233 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:

234 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:

235 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:

236 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:

237 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:

238 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:

239 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:

240 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:

241 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {

242 StringRef FeatureList(

243 getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));

245 CallerFeatureMap)) {

246 Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)

247 << FD->getDeclName() << FeatureList;

248 return false;

249 }

250

251 unsigned ArgCount = TheCall->getNumArgs() - 1;

253

254 return (SemaRef.BuiltinConstantArg(TheCall, 1, Result)) ||

255 (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||

256 (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));

257 }

258 default:

259 return false;

260 }

261

263 auto ArgExpr = Arg.get();

265

266 if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext()))

267 return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int)

268 << ArgExpr->getType();

269 auto Ord = ArgResult.Val.getInt().getZExtValue();

270

271

272

273 if (!llvm::isValidAtomicOrderingCABI(Ord))

274 return Diag(ArgExpr->getBeginLoc(),

275 diag::warn_atomic_op_has_invalid_memory_order)

276 << 0 << ArgExpr->getSourceRange();

277 switch (static_castllvm::AtomicOrderingCABI\(Ord)) {

278 case llvm::AtomicOrderingCABI::relaxed:

279 case llvm::AtomicOrderingCABI::consume:

280 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)

281 return Diag(ArgExpr->getBeginLoc(),

282 diag::warn_atomic_op_has_invalid_memory_order)

283 << 0 << ArgExpr->getSourceRange();

284 break;

285 case llvm::AtomicOrderingCABI::acquire:

286 case llvm::AtomicOrderingCABI::release:

287 case llvm::AtomicOrderingCABI::acq_rel:

288 case llvm::AtomicOrderingCABI::seq_cst:

289 break;

290 }

291

292 Arg = TheCall->getArg(ScopeIndex);

293 ArgExpr = Arg.get();

295

296 if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext()))

297 return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal)

298 << ArgExpr->getType();

299

300 return false;

301}

304 bool Fail = false;

305

306

310 if (AS != llvm::AMDGPUAS::FLAT_ADDRESS &&

311 AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) {

312 Fail = true;

313 Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as)

315 }

316

317

318 Expr *AtomicOrdArg = TheCall->getArg(IsStore ? 2 : 1);

321 llvm_unreachable("Intrinsic requires imm for atomic ordering argument!");

322 auto Ord =

323 llvm::AtomicOrderingCABI(AtomicOrdArgRes.Val.getInt().getZExtValue());

324

325

326

327 if (!llvm::isValidAtomicOrderingCABI((unsigned)Ord) ||

328 (Ord == llvm::AtomicOrderingCABI::acq_rel) ||

329 Ord == (IsStore ? llvm::AtomicOrderingCABI::acquire

330 : llvm::AtomicOrderingCABI::release)) {

332 diag::warn_atomic_op_has_invalid_memory_order)

334 }

335

336

339 Fail = true;

340 Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)

342 }

343

344 return Fail;

345}

348 unsigned NumDataArgs) {

349 assert(NumDataArgs <= 2);

350 if (SemaRef.checkArgCountRange(TheCall, NumArgs, NumArgs))

351 return true;

352 Expr *Args[2];

354 for (unsigned I = 0; I != NumDataArgs; ++I) {

355 Args[I] = TheCall->getArg(I);

356 ArgTys[I] = Args[I]->getType();

357

358 if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {

359 SemaRef.Diag(Args[I]->getBeginLoc(),

360 diag::err_typecheck_cond_expect_int_float)

362 return true;

363 }

364 }

365 if (NumDataArgs < 2)

366 return false;

367

368 if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1]))

369 return false;

370

371 if (((ArgTys[0]->isUnsignedIntegerType() &&

372 ArgTys[1]->isSignedIntegerType()) ||

373 (ArgTys[0]->isSignedIntegerType() &&

374 ArgTys[1]->isUnsignedIntegerType())) &&

377 return false;

378

379 SemaRef.Diag(Args[1]->getBeginLoc(),

380 diag::err_typecheck_call_different_arg_types)

381 << ArgTys[0] << ArgTys[1];

382 return true;

383}

387 const AMDGPUFlatWorkGroupSizeAttr &Attr) {

388

389

391 return false;

392

393 uint32_t Min = 0;

395 return true;

396

397 uint32_t Max = 0;

399 return true;

400

401 if (Min == 0 && Max != 0) {

403 << &Attr << 0;

404 return true;

405 }

408 << &Attr << 1;

409 return true;

410 }

411

412 return false;

413}

417 Expr *MinExpr, Expr *MaxExpr) {

419 AMDGPUFlatWorkGroupSizeAttr TmpAttr(Context, CI, MinExpr, MaxExpr);

420

422 return nullptr;

423 return ::new (Context)

424 AMDGPUFlatWorkGroupSizeAttr(Context, CI, MinExpr, MaxExpr);

425}

429 Expr *MinExpr, Expr *MaxExpr) {

432}

443 Expr *MaxExpr,

444 const AMDGPUWavesPerEUAttr &Attr) {

447 return true;

448

449

450

452 return false;

453

454 uint32_t Min = 0;

456 return true;

457

458 uint32_t Max = 0;

460 return true;

461

462 if (Min == 0 && Max != 0) {

464 << &Attr << 0;

465 return true;

466 }

469 << &Attr << 1;

470 return true;

471 }

472

473 return false;

474}

478 Expr *MinExpr, Expr *MaxExpr) {

480 AMDGPUWavesPerEUAttr TmpAttr(Context, CI, MinExpr, MaxExpr);

481

483 return nullptr;

484

485 return ::new (Context) AMDGPUWavesPerEUAttr(Context, CI, MinExpr, MaxExpr);

486}

489 Expr *MinExpr, Expr *MaxExpr) {

492}

505 uint32_t NumSGPR = 0;

507 if (SemaRef.checkUInt32Argument(AL, NumSGPRExpr, NumSGPR))

508 return;

509

511 AMDGPUNumSGPRAttr(getASTContext(), AL, NumSGPR));

512}

515 uint32_t NumVGPR = 0;

517 if (SemaRef.checkUInt32Argument(AL, NumVGPRExpr, NumVGPR))

518 return;

519

521 AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));

522}

527 const AMDGPUMaxNumWorkGroupsAttr &Attr) {

531 return true;

532

533

534

537 return false;

538

539 uint32_t NumWG = 0;

540 Expr *Exprs[3] = {XExpr, YExpr, ZExpr};

541 for (int i = 0; i < 3; i++) {

542 if (Exprs[i]) {

544 true))

545 return true;

546 if (NumWG == 0) {

547 S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)

549 return true;

550 }

551 }

552 }

553

554 return false;

555}

560 AMDGPUMaxNumWorkGroupsAttr TmpAttr(Context, CI, XExpr, YExpr, ZExpr);

561 assert(SemaRef.isSFINAEContext() &&

562 "Can't produce SFINAE diagnostic pointing to temporary attribute");

563

565 TmpAttr))

566 return nullptr;

567

568 return ::new (Context)

569 AMDGPUMaxNumWorkGroupsAttr(Context, CI, XExpr, YExpr, ZExpr);

570}