36#include "llvm/IR/IntrinsicsAArch64.h"
37#include "llvm/IR/IntrinsicsAMDGPU.h"
38#include "llvm/IR/IntrinsicsARM.h"
39#include "llvm/IR/IntrinsicsNVPTX.h"
40#include "llvm/IR/IntrinsicsRISCV.h"
41#include "llvm/IR/IntrinsicsWebAssembly.h"
42#include "llvm/IR/IntrinsicsX86.h"
65#define DEBUG_TYPE "auto-upgrade"
84 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
99 Type *LastArgType =
F->getFunctionType()->getParamType(
100 F->getFunctionType()->getNumParams() - 1);
115 if (
F->getReturnType()->isVectorTy())
128 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
129 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
146 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
147 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
161 if (
F->getReturnType()->getScalarType()->isBFloatTy())
171 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
185 if (Name.consume_front(
"avx."))
186 return (Name.starts_with(
"blend.p") ||
187 Name ==
"cvt.ps2.pd.256" ||
188 Name ==
"cvtdq2.pd.256" ||
189 Name ==
"cvtdq2.ps.256" ||
190 Name.starts_with(
"movnt.") ||
191 Name.starts_with(
"sqrt.p") ||
192 Name.starts_with(
"storeu.") ||
193 Name.starts_with(
"vbroadcast.s") ||
194 Name.starts_with(
"vbroadcastf128") ||
195 Name.starts_with(
"vextractf128.") ||
196 Name.starts_with(
"vinsertf128.") ||
197 Name.starts_with(
"vperm2f128.") ||
198 Name.starts_with(
"vpermil."));
200 if (Name.consume_front(
"avx2."))
201 return (Name ==
"movntdqa" ||
202 Name.starts_with(
"pabs.") ||
203 Name.starts_with(
"padds.") ||
204 Name.starts_with(
"paddus.") ||
205 Name.starts_with(
"pblendd.") ||
207 Name.starts_with(
"pbroadcast") ||
208 Name.starts_with(
"pcmpeq.") ||
209 Name.starts_with(
"pcmpgt.") ||
210 Name.starts_with(
"pmax") ||
211 Name.starts_with(
"pmin") ||
212 Name.starts_with(
"pmovsx") ||
213 Name.starts_with(
"pmovzx") ||
214 Name.starts_with(
"pmulh.w") ||
215 Name.starts_with(
"pmulhu.w") ||
217 Name ==
"pmulu.dq" ||
218 Name.starts_with(
"psll.dq") ||
219 Name.starts_with(
"psrl.dq") ||
220 Name.starts_with(
"psubs.") ||
221 Name.starts_with(
"psubus.") ||
222 Name.starts_with(
"vbroadcast") ||
223 Name ==
"vbroadcasti128" ||
224 Name ==
"vextracti128" ||
225 Name ==
"vinserti128" ||
226 Name ==
"vperm2i128");
228 if (Name.consume_front(
"avx512.")) {
229 if (Name.consume_front(
"mask."))
231 return (Name.starts_with(
"add.p") ||
232 Name.starts_with(
"and.") ||
233 Name.starts_with(
"andn.") ||
234 Name.starts_with(
"broadcast.s") ||
235 Name.starts_with(
"broadcastf32x4.") ||
236 Name.starts_with(
"broadcastf32x8.") ||
237 Name.starts_with(
"broadcastf64x2.") ||
238 Name.starts_with(
"broadcastf64x4.") ||
239 Name.starts_with(
"broadcasti32x4.") ||
240 Name.starts_with(
"broadcasti32x8.") ||
241 Name.starts_with(
"broadcasti64x2.") ||
242 Name.starts_with(
"broadcasti64x4.") ||
243 Name.starts_with(
"cmp.b") ||
244 Name.starts_with(
"cmp.d") ||
245 Name.starts_with(
"cmp.q") ||
246 Name.starts_with(
"cmp.w") ||
247 Name.starts_with(
"compress.b") ||
248 Name.starts_with(
"compress.d") ||
249 Name.starts_with(
"compress.p") ||
250 Name.starts_with(
"compress.q") ||
251 Name.starts_with(
"compress.store.") ||
252 Name.starts_with(
"compress.w") ||
253 Name.starts_with(
"conflict.") ||
254 Name.starts_with(
"cvtdq2pd.") ||
255 Name.starts_with(
"cvtdq2ps.") ||
256 Name ==
"cvtpd2dq.256" ||
257 Name ==
"cvtpd2ps.256" ||
258 Name ==
"cvtps2pd.128" ||
259 Name ==
"cvtps2pd.256" ||
260 Name.starts_with(
"cvtqq2pd.") ||
261 Name ==
"cvtqq2ps.256" ||
262 Name ==
"cvtqq2ps.512" ||
263 Name ==
"cvttpd2dq.256" ||
264 Name ==
"cvttps2dq.128" ||
265 Name ==
"cvttps2dq.256" ||
266 Name.starts_with(
"cvtudq2pd.") ||
267 Name.starts_with(
"cvtudq2ps.") ||
268 Name.starts_with(
"cvtuqq2pd.") ||
269 Name ==
"cvtuqq2ps.256" ||
270 Name ==
"cvtuqq2ps.512" ||
271 Name.starts_with(
"dbpsadbw.") ||
272 Name.starts_with(
"div.p") ||
273 Name.starts_with(
"expand.b") ||
274 Name.starts_with(
"expand.d") ||
275 Name.starts_with(
"expand.load.") ||
276 Name.starts_with(
"expand.p") ||
277 Name.starts_with(
"expand.q") ||
278 Name.starts_with(
"expand.w") ||
279 Name.starts_with(
"fpclass.p") ||
280 Name.starts_with(
"insert") ||
281 Name.starts_with(
"load.") ||
282 Name.starts_with(
"loadu.") ||
283 Name.starts_with(
"lzcnt.") ||
284 Name.starts_with(
"max.p") ||
285 Name.starts_with(
"min.p") ||
286 Name.starts_with(
"movddup") ||
287 Name.starts_with(
"move.s") ||
288 Name.starts_with(
"movshdup") ||
289 Name.starts_with(
"movsldup") ||
290 Name.starts_with(
"mul.p") ||
291 Name.starts_with(
"or.") ||
292 Name.starts_with(
"pabs.") ||
293 Name.starts_with(
"packssdw.") ||
294 Name.starts_with(
"packsswb.") ||
295 Name.starts_with(
"packusdw.") ||
296 Name.starts_with(
"packuswb.") ||
297 Name.starts_with(
"padd.") ||
298 Name.starts_with(
"padds.") ||
299 Name.starts_with(
"paddus.") ||
300 Name.starts_with(
"palignr.") ||
301 Name.starts_with(
"pand.") ||
302 Name.starts_with(
"pandn.") ||
303 Name.starts_with(
"pavg") ||
304 Name.starts_with(
"pbroadcast") ||
305 Name.starts_with(
"pcmpeq.") ||
306 Name.starts_with(
"pcmpgt.") ||
307 Name.starts_with(
"perm.df.") ||
308 Name.starts_with(
"perm.di.") ||
309 Name.starts_with(
"permvar.") ||
310 Name.starts_with(
"pmaddubs.w.") ||
311 Name.starts_with(
"pmaddw.d.") ||
312 Name.starts_with(
"pmax") ||
313 Name.starts_with(
"pmin") ||
314 Name ==
"pmov.qd.256" ||
315 Name ==
"pmov.qd.512" ||
316 Name ==
"pmov.wb.256" ||
317 Name ==
"pmov.wb.512" ||
318 Name.starts_with(
"pmovsx") ||
319 Name.starts_with(
"pmovzx") ||
320 Name.starts_with(
"pmul.dq.") ||
321 Name.starts_with(
"pmul.hr.sw.") ||
322 Name.starts_with(
"pmulh.w.") ||
323 Name.starts_with(
"pmulhu.w.") ||
324 Name.starts_with(
"pmull.") ||
325 Name.starts_with(
"pmultishift.qb.") ||
326 Name.starts_with(
"pmulu.dq.") ||
327 Name.starts_with(
"por.") ||
328 Name.starts_with(
"prol.") ||
329 Name.starts_with(
"prolv.") ||
330 Name.starts_with(
"pror.") ||
331 Name.starts_with(
"prorv.") ||
332 Name.starts_with(
"pshuf.b.") ||
333 Name.starts_with(
"pshuf.d.") ||
334 Name.starts_with(
"pshufh.w.") ||
335 Name.starts_with(
"pshufl.w.") ||
336 Name.starts_with(
"psll.d") ||
337 Name.starts_with(
"psll.q") ||
338 Name.starts_with(
"psll.w") ||
339 Name.starts_with(
"pslli") ||
340 Name.starts_with(
"psllv") ||
341 Name.starts_with(
"psra.d") ||
342 Name.starts_with(
"psra.q") ||
343 Name.starts_with(
"psra.w") ||
344 Name.starts_with(
"psrai") ||
345 Name.starts_with(
"psrav") ||
346 Name.starts_with(
"psrl.d") ||
347 Name.starts_with(
"psrl.q") ||
348 Name.starts_with(
"psrl.w") ||
349 Name.starts_with(
"psrli") ||
350 Name.starts_with(
"psrlv") ||
351 Name.starts_with(
"psub.") ||
352 Name.starts_with(
"psubs.") ||
353 Name.starts_with(
"psubus.") ||
354 Name.starts_with(
"pternlog.") ||
355 Name.starts_with(
"punpckh") ||
356 Name.starts_with(
"punpckl") ||
357 Name.starts_with(
"pxor.") ||
358 Name.starts_with(
"shuf.f") ||
359 Name.starts_with(
"shuf.i") ||
360 Name.starts_with(
"shuf.p") ||
361 Name.starts_with(
"sqrt.p") ||
362 Name.starts_with(
"store.b.") ||
363 Name.starts_with(
"store.d.") ||
364 Name.starts_with(
"store.p") ||
365 Name.starts_with(
"store.q.") ||
366 Name.starts_with(
"store.w.") ||
367 Name ==
"store.ss" ||
368 Name.starts_with(
"storeu.") ||
369 Name.starts_with(
"sub.p") ||
370 Name.starts_with(
"ucmp.") ||
371 Name.starts_with(
"unpckh.") ||
372 Name.starts_with(
"unpckl.") ||
373 Name.starts_with(
"valign.") ||
374 Name ==
"vcvtph2ps.128" ||
375 Name ==
"vcvtph2ps.256" ||
376 Name.starts_with(
"vextract") ||
377 Name.starts_with(
"vfmadd.") ||
378 Name.starts_with(
"vfmaddsub.") ||
379 Name.starts_with(
"vfnmadd.") ||
380 Name.starts_with(
"vfnmsub.") ||
381 Name.starts_with(
"vpdpbusd.") ||
382 Name.starts_with(
"vpdpbusds.") ||
383 Name.starts_with(
"vpdpwssd.") ||
384 Name.starts_with(
"vpdpwssds.") ||
385 Name.starts_with(
"vpermi2var.") ||
386 Name.starts_with(
"vpermil.p") ||
387 Name.starts_with(
"vpermilvar.") ||
388 Name.starts_with(
"vpermt2var.") ||
389 Name.starts_with(
"vpmadd52") ||
390 Name.starts_with(
"vpshld.") ||
391 Name.starts_with(
"vpshldv.") ||
392 Name.starts_with(
"vpshrd.") ||
393 Name.starts_with(
"vpshrdv.") ||
394 Name.starts_with(
"vpshufbitqmb.") ||
395 Name.starts_with(
"xor."));
397 if (Name.consume_front(
"mask3."))
399 return (Name.starts_with(
"vfmadd.") ||
400 Name.starts_with(
"vfmaddsub.") ||
401 Name.starts_with(
"vfmsub.") ||
402 Name.starts_with(
"vfmsubadd.") ||
403 Name.starts_with(
"vfnmsub."));
405 if (Name.consume_front(
"maskz."))
407 return (Name.starts_with(
"pternlog.") ||
408 Name.starts_with(
"vfmadd.") ||
409 Name.starts_with(
"vfmaddsub.") ||
410 Name.starts_with(
"vpdpbusd.") ||
411 Name.starts_with(
"vpdpbusds.") ||
412 Name.starts_with(
"vpdpwssd.") ||
413 Name.starts_with(
"vpdpwssds.") ||
414 Name.starts_with(
"vpermt2var.") ||
415 Name.starts_with(
"vpmadd52") ||
416 Name.starts_with(
"vpshldv.") ||
417 Name.starts_with(
"vpshrdv."));
420 return (Name ==
"movntdqa" ||
421 Name ==
"pmul.dq.512" ||
422 Name ==
"pmulu.dq.512" ||
423 Name.starts_with(
"broadcastm") ||
424 Name.starts_with(
"cmp.p") ||
425 Name.starts_with(
"cvtb2mask.") ||
426 Name.starts_with(
"cvtd2mask.") ||
427 Name.starts_with(
"cvtmask2") ||
428 Name.starts_with(
"cvtq2mask.") ||
429 Name ==
"cvtusi2sd" ||
430 Name.starts_with(
"cvtw2mask.") ||
435 Name ==
"kortestc.w" ||
436 Name ==
"kortestz.w" ||
437 Name.starts_with(
"kunpck") ||
440 Name.starts_with(
"padds.") ||
441 Name.starts_with(
"pbroadcast") ||
442 Name.starts_with(
"pmulh.w") ||
443 Name.starts_with(
"pmulhu.w") ||
444 Name.starts_with(
"prol") ||
445 Name.starts_with(
"pror") ||
446 Name.starts_with(
"psll.dq") ||
447 Name.starts_with(
"psrl.dq") ||
448 Name.starts_with(
"psubs.") ||
449 Name.starts_with(
"ptestm") ||
450 Name.starts_with(
"ptestnm") ||
451 Name.starts_with(
"storent.") ||
452 Name.starts_with(
"vbroadcast.s") ||
453 Name.starts_with(
"vpshld.") ||
454 Name.starts_with(
"vpshrd."));
457 if (Name.consume_front(
"fma."))
458 return (Name.starts_with(
"vfmadd.") ||
459 Name.starts_with(
"vfmsub.") ||
460 Name.starts_with(
"vfmsubadd.") ||
461 Name.starts_with(
"vfnmadd.") ||
462 Name.starts_with(
"vfnmsub."));
464 if (Name.consume_front(
"fma4."))
465 return Name.starts_with(
"vfmadd.s");
467 if (Name.consume_front(
"sse."))
468 return (Name ==
"add.ss" ||
469 Name ==
"cvtsi2ss" ||
470 Name ==
"cvtsi642ss" ||
473 Name.starts_with(
"sqrt.p") ||
475 Name.starts_with(
"storeu.") ||
478 if (Name.consume_front(
"sse2."))
479 return (Name ==
"add.sd" ||
480 Name ==
"cvtdq2pd" ||
481 Name ==
"cvtdq2ps" ||
482 Name ==
"cvtps2pd" ||
483 Name ==
"cvtsi2sd" ||
484 Name ==
"cvtsi642sd" ||
485 Name ==
"cvtss2sd" ||
488 Name.starts_with(
"padds.") ||
489 Name.starts_with(
"paddus.") ||
490 Name.starts_with(
"pcmpeq.") ||
491 Name.starts_with(
"pcmpgt.") ||
497 Name ==
"pmulhu.w" ||
498 Name ==
"pmulu.dq" ||
499 Name.starts_with(
"pshuf") ||
500 Name.starts_with(
"psll.dq") ||
501 Name.starts_with(
"psrl.dq") ||
502 Name.starts_with(
"psubs.") ||
503 Name.starts_with(
"psubus.") ||
504 Name.starts_with(
"sqrt.p") ||
506 Name ==
"storel.dq" ||
507 Name.starts_with(
"storeu.") ||
510 if (Name.consume_front(
"sse41."))
511 return (Name.starts_with(
"blendp") ||
512 Name ==
"movntdqa" ||
522 Name.starts_with(
"pmovsx") ||
523 Name.starts_with(
"pmovzx") ||
526 if (Name.consume_front(
"sse42."))
527 return Name ==
"crc32.64.8";
529 if (Name.consume_front(
"sse4a."))
530 return Name.starts_with(
"movnt.");
532 if (Name.consume_front(
"ssse3."))
533 return (Name ==
"pabs.b.128" ||
534 Name ==
"pabs.d.128" ||
535 Name ==
"pabs.w.128");
537 if (Name.consume_front(
"xop."))
538 return (Name ==
"vpcmov" ||
539 Name ==
"vpcmov.256" ||
540 Name.starts_with(
"vpcom") ||
541 Name.starts_with(
"vprot"));
543 if (Name.consume_front(
"bmi."))
544 return (Name.starts_with(
"pdep.") ||
545 Name.starts_with(
"pext."));
547 return (Name ==
"addcarry.u32" ||
548 Name ==
"addcarry.u64" ||
549 Name ==
"addcarryx.u32" ||
550 Name ==
"addcarryx.u64" ||
551 Name ==
"subborrow.u32" ||
552 Name ==
"subborrow.u64" ||
553 Name.starts_with(
"vcvtph2ps."));
559 if (!Name.consume_front(
"x86."))
567 if (Name ==
"rdtscp") {
569 if (
F->getFunctionType()->getNumParams() == 0)
574 Intrinsic::x86_rdtscp);
581 if (Name.consume_front(
"sse41.ptest")) {
583 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
584 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
585 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
598 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
599 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
600 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
601 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
602 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
603 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
608 if (Name.consume_front(
"avx512.")) {
609 if (Name.consume_front(
"mask.cmp.")) {
612 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
613 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
614 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
615 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
616 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
617 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
621 }
else if (Name.starts_with(
"vpdpbusd.") ||
622 Name.starts_with(
"vpdpbusds.")) {
625 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
626 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
627 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
628 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
629 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
630 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
634 }
else if (Name.starts_with(
"vpdpwssd.") ||
635 Name.starts_with(
"vpdpwssds.")) {
638 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
639 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
640 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
641 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
642 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
643 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
651 if (Name.consume_front(
"avx2.")) {
652 if (Name.consume_front(
"vpdpb")) {
655 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
656 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
657 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
658 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
659 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
660 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
661 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
662 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
663 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
664 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
665 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
666 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
670 }
else if (Name.consume_front(
"vpdpw")) {
673 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
674 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
675 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
676 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
677 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
678 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
679 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
680 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
681 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
682 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
683 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
684 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
692 if (Name.consume_front(
"avx10.")) {
693 if (Name.consume_front(
"vpdpb")) {
696 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
697 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
698 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
699 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
700 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
701 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
705 }
else if (Name.consume_front(
"vpdpw")) {
707 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
708 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
709 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
710 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
711 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
712 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
720 if (Name.consume_front(
"avx512bf16.")) {
723 .
Case(
"cvtne2ps2bf16.128",
724 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
725 .
Case(
"cvtne2ps2bf16.256",
726 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
727 .
Case(
"cvtne2ps2bf16.512",
728 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
729 .
Case(
"mask.cvtneps2bf16.128",
730 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
731 .
Case(
"cvtneps2bf16.256",
732 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
733 .
Case(
"cvtneps2bf16.512",
734 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
741 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
742 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
743 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
750 if (Name.consume_front(
"xop.")) {
752 if (Name.starts_with(
"vpermil2")) {
755 auto Idx =
F->getFunctionType()->getParamType(2);
756 if (Idx->isFPOrFPVectorTy()) {
757 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
758 unsigned EltSize = Idx->getScalarSizeInBits();
759 if (EltSize == 64 && IdxSize == 128)
760 ID = Intrinsic::x86_xop_vpermil2pd;
761 else if (EltSize == 32 && IdxSize == 128)
762 ID = Intrinsic::x86_xop_vpermil2ps;
763 else if (EltSize == 64 && IdxSize == 256)
764 ID = Intrinsic::x86_xop_vpermil2pd_256;
766 ID = Intrinsic::x86_xop_vpermil2ps_256;
768 }
else if (
F->arg_size() == 2)
771 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
772 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
783 if (Name ==
"seh.recoverfp") {
785 Intrinsic::eh_recoverfp);
797 if (Name.starts_with(
"rbit")) {
800 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
804 if (Name ==
"thread.pointer") {
807 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
811 bool Neon = Name.consume_front(
"neon.");
816 if (Name.consume_front(
"bfdot.")) {
820 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
825 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
826 assert((OperandWidth == 64 || OperandWidth == 128) &&
827 "Unexpected operand width");
829 std::array<Type *, 2> Tys{
840 if (Name.consume_front(
"bfm")) {
842 if (Name.consume_back(
".v4f32.v16i8")) {
888 F->arg_begin()->getType());
892 if (Name.consume_front(
"vst")) {
894 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
898 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
899 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
902 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
903 Intrinsic::arm_neon_vst4lane};
905 auto fArgs =
F->getFunctionType()->params();
906 Type *Tys[] = {fArgs[0], fArgs[1]};
909 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
912 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
921 if (Name.consume_front(
"mve.")) {
923 if (Name ==
"vctp64") {
933 if (Name.starts_with(
"vrintn.v")) {
935 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
940 if (Name.consume_back(
".v4i1")) {
942 if (Name.consume_back(
".predicated.v2i64.v4i32"))
944 return Name ==
"mull.int" || Name ==
"vqdmull";
946 if (Name.consume_back(
".v2i64")) {
948 bool IsGather = Name.consume_front(
"vldr.gather.");
949 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
950 if (Name.consume_front(
"base.")) {
952 Name.consume_front(
"wb.");
955 return Name ==
"predicated.v2i64";
958 if (Name.consume_front(
"offset.predicated."))
959 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
960 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
973 if (Name.consume_front(
"cde.vcx")) {
975 if (Name.consume_back(
".predicated.v2i64.v4i1"))
977 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
978 Name ==
"3q" || Name ==
"3qa";
992 F->arg_begin()->getType());
998 .
Case(
"smax", Intrinsic::smax)
999 .
Case(
"smin", Intrinsic::smin)
1000 .
Case(
"umax", Intrinsic::umax)
1001 .
Case(
"umin", Intrinsic::umin)
1004 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1007 F->getReturnType());
1011 if (Name.starts_with(
"addp")) {
1013 if (
F->arg_size() != 2)
1016 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1018 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1024 if (Name.starts_with(
"bfcvt")) {
1030 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1037 if (Name.consume_front(
"sve.")) {
1039 if (Name.consume_front(
"bf")) {
1040 if (Name ==
"mmla") {
1041 Type *Tys[] = {
F->getReturnType(),
1042 std::next(
F->arg_begin())->getType()};
1044 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1047 if (Name.consume_back(
".lane")) {
1051 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1052 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1053 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1065 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1070 if (Name.consume_front(
"convert.from.svbool")) {
1073 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1076 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1081 if (Name.consume_front(
"convert.to.svbool")) {
1084 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1087 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1092 if (Name.consume_front(
"addqv")) {
1094 if (!
F->getReturnType()->isFPOrFPVectorTy())
1097 auto Args =
F->getFunctionType()->params();
1098 Type *Tys[] = {
F->getReturnType(), Args[1]};
1100 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1104 if (Name.consume_front(
"ld")) {
1106 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1107 if (LdRegex.
match(Name)) {
1113 "Expected 2 arguments for ld* intrinsic.");
1114 Type *PtrTy =
F->getArg(1)->getType();
1117 Intrinsic::aarch64_sve_ld2_sret,
1118 Intrinsic::aarch64_sve_ld3_sret,
1119 Intrinsic::aarch64_sve_ld4_sret,
1122 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1128 if (Name.consume_front(
"tuple.")) {
1130 if (Name.starts_with(
"get")) {
1132 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1134 F->getParent(), Intrinsic::vector_extract, Tys);
1138 if (Name.starts_with(
"set")) {
1140 auto Args =
F->getFunctionType()->params();
1141 Type *Tys[] = {Args[0], Args[2], Args[1]};
1143 F->getParent(), Intrinsic::vector_insert, Tys);
1147 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1148 if (CreateTupleRegex.
match(Name)) {
1150 auto Args =
F->getFunctionType()->params();
1151 Type *Tys[] = {
F->getReturnType(), Args[1]};
1153 F->getParent(), Intrinsic::vector_insert, Tys);
1159 if (Name.starts_with(
"rev.nxv")) {
1162 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1168 if (Name.consume_front(
"sme.")) {
1170 if (Name.consume_front(
"ftmopa.")) {
1175 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1176 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1194#define NVVM_TMA_G2S_MODES(M) \
1195 M(tile_1d, "tile.1d") \
1196 M(tile_2d, "tile.2d") \
1197 M(tile_3d, "tile.3d") \
1198 M(tile_4d, "tile.4d") \
1199 M(tile_5d, "tile.5d") \
1200 M(tile_gather4_2d, "tile.gather4.2d") \
1201 M(im2col_3d, "im2col.3d") \
1202 M(im2col_4d, "im2col.4d") \
1203 M(im2col_5d, "im2col.5d") \
1204 M(im2col_w_3d, "im2col.w.3d") \
1205 M(im2col_w_4d, "im2col.w.4d") \
1206 M(im2col_w_5d, "im2col.w.5d") \
1207 M(im2col_w_128_3d, "im2col.w.128.3d") \
1208 M(im2col_w_128_4d, "im2col.w.128.4d") \
1209 M(im2col_w_128_5d, "im2col.w.128.5d")
1221 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1224#define G2S_ID(ID_SUFFIX, NAME) \
1225 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1235 size_t NumParams =
F->getFunctionType()->getNumParams();
1239 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1246 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1247 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1248 "expected the i64 cache-hint after the multicast mask");
1249 Type *MaskTy = Params[MaskIdx];
1264 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1267#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1268 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1280 if (!
F->getFunctionType()
1281 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1297 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1302 size_t NumParams =
F->getFunctionType()->getNumParams();
1303 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1307 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1312 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1325 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1330 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1333 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1353 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1356 auto [RedOpName, ShapeName] = Name.split(
'.');
1361 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1362 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1363 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1364 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1365 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1366 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1367 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1368 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1374 if (!Name.consume_front(
"mbarrier.") ||
1375 !Name.consume_back(
".space.cluster") ||
F->arg_size() != 2)
1388 if (Name.consume_front(
"mapa.shared.cluster"))
1389 if (
F->getReturnType()->getPointerAddressSpace() ==
1391 return Intrinsic::nvvm_mapa_shared_cluster;
1393 if (Name.consume_front(
"cp.async.bulk.")) {
1396 .
Case(
"shared.cta.to.cluster",
1397 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1401 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1411 if (!Name.consume_front(
"tcgen05.commit."))
1414 if (Name.consume_front(
"shared."))
1416 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1417 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1420 if (Name.consume_front(
"mc.shared.")) {
1422 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1426 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1427 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1436 if (
F->arg_size() != 2)
1440 .
Case(
"tcgen05.alloc.shared.cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1441 .
Case(
"tcgen05.alloc.shared.cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1446 if (Name.consume_front(
"fmax."))
1448 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1449 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1450 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1451 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1452 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1453 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1454 .
Case(
"ftz.nan.xorsign.abs.bf16",
1455 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1456 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1457 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1458 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1459 .
Case(
"ftz.xorsign.abs.bf16x2",
1460 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1461 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1462 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1463 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1464 .
Case(
"nan.xorsign.abs.bf16x2",
1465 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1466 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1467 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1470 if (Name.consume_front(
"fmin."))
1472 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1473 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1474 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1475 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1476 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1477 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1478 .
Case(
"ftz.nan.xorsign.abs.bf16",
1479 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1480 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1481 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1482 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1483 .
Case(
"ftz.xorsign.abs.bf16x2",
1484 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1485 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1486 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1487 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1488 .
Case(
"nan.xorsign.abs.bf16x2",
1489 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1490 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1491 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1494 if (Name.consume_front(
"neg."))
1496 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1497 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1506 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1511 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1514 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1517 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1518 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1528 {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat,
1532 {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat,
1535 {Intrinsic::nvvm_fma, Intrinsic::nvvm_fma_sat, Intrinsic::nvvm_fma_relu},
1536 {Intrinsic::nvvm_fma_ftz, Intrinsic::nvvm_fma_ftz_sat,
1537 Intrinsic::nvvm_fma_ftz_relu}};
1543 bool IsFTZ,
bool IsSat,
1548 .
Case(
"f", !IsRelu && !(IsSat && IsFMul))
1549 .
Case(
"d", !IsFTZ && !IsSat && !IsRelu)
1550 .
Case(
"f16", IsRN && (IsFMA || IsSat))
1551 .
Case(
"v2f16", IsRN && !IsFMA && IsSat)
1552 .
Case(
"f16x2", IsRN && IsFMA)
1553 .
Cases({
"bf16",
"bf16x2"}, IsRN && IsFMA && !IsFTZ && !IsSat)
1557static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1559 auto [Modifiers,
Type] = Name.rsplit(
'.');
1569 return std::nullopt;
1571 StringRef Rest = Modifiers.drop_front(2);
1576 return std::nullopt;
1578 const Intrinsic::ID IID = IIDs[IsFTZ][IsSat ? 1 : IsRelu ? 2 : 0];
1580 return std::nullopt;
1586 return std::nullopt;
1592 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1595 return Intrinsic::nvvm_mbarrier_init;
1599 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1600 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1601 Name.consume_front(
"param");
1605 if (!Name.consume_front(
"vp."))
1634 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1635 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1642 if (!Name.consume_front(
"vp."))
1662 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1663 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1668 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1680 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1691 if (Name.starts_with(
"to.fp16")) {
1695 FuncTy->getReturnType());
1698 if (Name.starts_with(
"from.fp16")) {
1702 FuncTy->getReturnType());
1712 if (Defaults.empty())
1715 unsigned FullArgCount = FirstDefault + Defaults.size();
1718 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1721 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1723 NumMissingTrailingParams))
1726 return FullArgCount;
1733 unsigned FullArgCount =
1735 if (FullArgCount == 0)
1741 "total number of default args does not match intrinsic signature");
1746 bool CanUpgradeDebugIntrinsicsToRecords) {
1747 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1752 if (!Name.consume_front(
"llvm.") || Name.empty())
1758 bool IsArm = Name.consume_front(
"arm.");
1759 if (IsArm || Name.consume_front(
"aarch64.")) {
1765 if (Name.consume_front(
"amdgcn.")) {
1766 if (Name ==
"alignbit") {
1769 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1773 if (Name.consume_front(
"atomic.")) {
1774 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1775 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1784 if (Name.starts_with(
"addrspacecast.nonnull")) {
1791 switch (
F->getIntrinsicID()) {
1795 case Intrinsic::amdgcn_asyncmark:
1796 case Intrinsic::amdgcn_wait_asyncmark: {
1798 if (
F->arg_size() == (ID == Intrinsic::amdgcn_asyncmark ? 0u : 1u)) {
1806 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1807 if (
F->arg_size() == 7) {
1812 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1813 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1814 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1815 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1816 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1817 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1818 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1819 if (
F->arg_size() == 8) {
1826 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1827 Name.consume_front(
"flat.atomic.")) {
1828 if (Name.starts_with(
"fadd") ||
1830 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1831 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1839 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1844 if (Name.starts_with(
"ldexp.")) {
1847 F->getParent(), Intrinsic::ldexp,
1848 {F->getReturnType(), F->getArg(1)->getType()});
1857 if (
F->arg_size() == 1) {
1858 if (Name.consume_front(
"convert.")) {
1872 F->arg_begin()->getType());
1878 if (Name ==
"coro.end" &&
1879 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1880 CoroEndID = Intrinsic::coro_end;
1881 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1882 CoroEndID = Intrinsic::coro_end_async;
1893 if (Name.consume_front(
"dbg.")) {
1895 if (CanUpgradeDebugIntrinsicsToRecords) {
1896 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1897 Name ==
"declare" || Name ==
"label") {
1906 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1909 Intrinsic::dbg_value);
1916 if (Name.consume_front(
"experimental.vector.")) {
1922 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1923 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1924 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1925 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1926 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1928 Intrinsic::vector_partial_reduce_add)
1931 const auto *FT =
F->getFunctionType();
1933 if (ID == Intrinsic::vector_extract ||
1934 ID == Intrinsic::vector_interleave2)
1937 if (ID != Intrinsic::vector_interleave2)
1939 if (ID == Intrinsic::vector_insert ||
1940 ID == Intrinsic::vector_partial_reduce_add)
1948 if (Name.consume_front(
"reduce.")) {
1950 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1951 if (R.match(Name, &
Groups))
1953 .
Case(
"add", Intrinsic::vector_reduce_add)
1954 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1955 .
Case(
"and", Intrinsic::vector_reduce_and)
1956 .
Case(
"or", Intrinsic::vector_reduce_or)
1957 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1958 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1959 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1960 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1961 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1962 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1963 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1968 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1973 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1974 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1979 auto Args =
F->getFunctionType()->params();
1981 {Args[V2 ? 1 : 0]});
1987 if (Name.consume_front(
"splice"))
1991 if (Name.consume_front(
"experimental.stepvector.")) {
1995 F->getParent(), ID,
F->getFunctionType()->getReturnType());
2000 if (Name.starts_with(
"flt.rounds")) {
2003 Intrinsic::get_rounding);
2008 if (Name.starts_with(
"invariant.group.barrier")) {
2010 auto Args =
F->getFunctionType()->params();
2011 Type* ObjectPtr[1] = {Args[0]};
2014 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
2019 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
2020 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
2021 if (IsLifetimeStart || IsLifetimeEnd) {
2022 if (
F->arg_size() == 2) {
2023 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
2024 : Intrinsic::lifetime_end;
2029 F->getArg(1)->getType());
2031 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2051 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2052 .StartsWith(
"memmove.", Intrinsic::memmove)
2054 if (
F->arg_size() == 5) {
2058 F->getFunctionType()->params().slice(0, 3);
2064 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2067 const auto *FT =
F->getFunctionType();
2068 Type *ParamTypes[2] = {
2069 FT->getParamType(0),
2073 Intrinsic::memset, ParamTypes);
2079 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2080 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2081 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2082 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2084 if (MaskedID &&
F->arg_size() == 4) {
2086 if (MaskedID == Intrinsic::masked_load ||
2087 MaskedID == Intrinsic::masked_gather) {
2089 F->getParent(), MaskedID,
2090 {F->getReturnType(), F->getArg(0)->getType()});
2094 F->getParent(), MaskedID,
2095 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2101 if (Name.consume_front(
"nvvm.")) {
2103 if (
F->arg_size() == 1) {
2106 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2107 .Case(
"clz.i", Intrinsic::ctlz)
2108 .
Case(
"popc.i", Intrinsic::ctpop)
2112 {F->getReturnType()});
2115 }
else if (
F->arg_size() == 2) {
2118 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2119 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2120 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2121 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2122 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2123 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2127 {F->getReturnType()});
2173 F->getParent(), IID,
F->getReturnType(),
2174 F->getFunctionType()->params());
2183 {F->getArg(0)->getType()});
2229 F->getArg(0)->getType());
2237 bool Expand =
false;
2238 if (Name.consume_front(
"abs."))
2241 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2242 else if (Name.consume_front(
"fabs."))
2244 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2245 else if (Name.consume_front(
"add."))
2248 else if (Name.consume_front(
"mul."))
2251 else if (Name.consume_front(
"ex2.approx."))
2254 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2255 else if (Name.consume_front(
"fma."))
2258 else if (Name.consume_front(
"atomic.load."))
2267 else if (Name.consume_front(
"atomic."))
2282 else if (Name.consume_front(
"bitcast."))
2285 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2286 else if (Name.consume_front(
"rotate."))
2288 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2289 else if (Name.consume_front(
"ptr.gen.to."))
2292 else if (Name.consume_front(
"ptr."))
2295 else if (Name.consume_front(
"ldg.global."))
2297 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2298 Name.starts_with(
"p."));
2301 .
Case(
"barrier0",
true)
2302 .
Case(
"barrier.n",
true)
2303 .
Case(
"barrier.sync.cnt",
true)
2304 .
Case(
"barrier.sync",
true)
2305 .
Case(
"barrier",
true)
2306 .
Case(
"bar.sync",
true)
2307 .
Case(
"barrier0.popc",
true)
2308 .
Case(
"barrier0.and",
true)
2309 .
Case(
"barrier0.or",
true)
2310 .
Case(
"clz.ll",
true)
2311 .
Case(
"popc.ll",
true)
2313 .
Case(
"swap.lo.hi.b64",
true)
2314 .
Case(
"tanh.approx.f32",
true)
2326 if (Name.starts_with(
"objectsize.")) {
2327 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2328 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2331 Intrinsic::objectsize, Tys);
2338 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2341 F->getParent(), Intrinsic::ptr_annotation,
2342 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2348 if (Name.consume_front(
"riscv.")) {
2351 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2352 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2353 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2354 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2357 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2370 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2371 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2380 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2381 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2382 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2383 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2388 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2397 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2399 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2408 if (Name ==
"stackprotectorcheck") {
2412 if (Name.starts_with(
"strip.invariant.group")) {
2417 F->getParent(), Intrinsic::launder_invariant_group,
2418 F->getReturnType());
2424 if (Name ==
"thread.pointer") {
2426 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2432 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2435 F->getParent(), Intrinsic::var_annotation,
2436 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2439 if (Name.consume_front(
"vector.splice")) {
2440 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2450 if (Name.consume_front(
"wasm.")) {
2453 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2454 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2455 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2460 F->getReturnType());
2464 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2466 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2468 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2487 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2497 std::string
Name =
F->getName().str();
2500 Name,
F->getParent());
2511 if (Result != std::nullopt) {
2528 bool CanUpgradeDebugIntrinsicsToRecords) {
2548 GV->
getName() ==
"llvm.global_dtors")) ||
2562 unsigned N =
Init->getNumOperands();
2563 std::vector<Constant *> NewCtors(
N);
2564 for (
unsigned i = 0; i !=
N; ++i) {
2567 Ctor->getAggregateElement(1),
2581 unsigned NumElts = ResultTy->getNumElements() * 8;
2585 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2595 for (
unsigned l = 0; l != NumElts; l += 16)
2596 for (
unsigned i = 0; i != 16; ++i) {
2597 unsigned Idx = NumElts + i - Shift;
2599 Idx -= NumElts - 16;
2600 Idxs[l + i] = Idx + l;
2603 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2607 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2615 unsigned NumElts = ResultTy->getNumElements() * 8;
2619 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2629 for (
unsigned l = 0; l != NumElts; l += 16)
2630 for (
unsigned i = 0; i != 16; ++i) {
2631 unsigned Idx = i + Shift;
2633 Idx += NumElts - 16;
2634 Idxs[l + i] = Idx + l;
2637 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2641 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2649 Mask = Builder.CreateBitCast(Mask, MaskTy);
2655 for (
unsigned i = 0; i != NumElts; ++i)
2657 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2668 if (
C->isAllOnesValue())
2673 return Builder.CreateSelect(Mask, Op0, Op1);
2680 if (
C->isAllOnesValue())
2684 Mask->getType()->getIntegerBitWidth());
2685 Mask = Builder.CreateBitCast(Mask, MaskTy);
2686 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2687 return Builder.CreateSelect(Mask, Op0, Op1);
2700 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2701 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2706 ShiftVal &= (NumElts - 1);
2715 if (ShiftVal > 16) {
2723 for (
unsigned l = 0; l < NumElts; l += 16) {
2724 for (
unsigned i = 0; i != 16; ++i) {
2725 unsigned Idx = ShiftVal + i;
2726 if (!IsVALIGN && Idx >= 16)
2727 Idx += NumElts - 16;
2728 Indices[l + i] = Idx + l;
2733 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2739 bool ZeroMask,
bool IndexForm) {
2742 unsigned EltWidth = Ty->getScalarSizeInBits();
2743 bool IsFloat = Ty->isFPOrFPVectorTy();
2745 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2746 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2747 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2748 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2749 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2750 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2751 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2752 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2753 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2754 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2755 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2756 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2757 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2758 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2759 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2760 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2761 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2762 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2763 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2764 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2765 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2766 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2767 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2768 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2769 else if (VecWidth == 128 && EltWidth == 16)
2770 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2771 else if (VecWidth == 256 && EltWidth == 16)
2772 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2773 else if (VecWidth == 512 && EltWidth == 16)
2774 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2775 else if (VecWidth == 128 && EltWidth == 8)
2776 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2777 else if (VecWidth == 256 && EltWidth == 8)
2778 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2779 else if (VecWidth == 512 && EltWidth == 8)
2780 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2791 Value *
V = Builder.CreateIntrinsic(IID, Args);
2803 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2814 bool IsRotateRight) {
2824 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2825 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2828 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2829 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2874 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2879 bool IsShiftRight,
bool ZeroMask) {
2893 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2894 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2897 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2898 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2913 const Align Alignment =
2915 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2920 if (
C->isAllOnesValue())
2921 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2926 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2932 const Align Alignment =
2941 if (
C->isAllOnesValue())
2942 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2947 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2953 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2954 {Op0, Builder.getInt1(
false)});
2969 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2970 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2971 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2972 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2973 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2976 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2977 LHS = Builder.CreateAnd(
LHS, Mask);
2978 RHS = Builder.CreateAnd(
RHS, Mask);
2995 if (!
C || !
C->isAllOnesValue())
2996 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
3001 for (
unsigned i = 0; i != NumElts; ++i)
3003 for (
unsigned i = NumElts; i != 8; ++i)
3004 Indices[i] = NumElts + i % NumElts;
3005 Vec = Builder.CreateShuffleVector(Vec,
3009 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
3021 }
else if (
CC == 7) {
3057 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3058 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3060 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3061 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3070 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3076 Name = Name.substr(12);
3081 if (Name.starts_with(
"max.p")) {
3082 if (VecWidth == 128 && EltWidth == 32)
3083 IID = Intrinsic::x86_sse_max_ps;
3084 else if (VecWidth == 128 && EltWidth == 64)
3085 IID = Intrinsic::x86_sse2_max_pd;
3086 else if (VecWidth == 256 && EltWidth == 32)
3087 IID = Intrinsic::x86_avx_max_ps_256;
3088 else if (VecWidth == 256 && EltWidth == 64)
3089 IID = Intrinsic::x86_avx_max_pd_256;
3092 }
else if (Name.starts_with(
"min.p")) {
3093 if (VecWidth == 128 && EltWidth == 32)
3094 IID = Intrinsic::x86_sse_min_ps;
3095 else if (VecWidth == 128 && EltWidth == 64)
3096 IID = Intrinsic::x86_sse2_min_pd;
3097 else if (VecWidth == 256 && EltWidth == 32)
3098 IID = Intrinsic::x86_avx_min_ps_256;
3099 else if (VecWidth == 256 && EltWidth == 64)
3100 IID = Intrinsic::x86_avx_min_pd_256;
3103 }
else if (Name.starts_with(
"pshuf.b.")) {
3104 if (VecWidth == 128)
3105 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3106 else if (VecWidth == 256)
3107 IID = Intrinsic::x86_avx2_pshuf_b;
3108 else if (VecWidth == 512)
3109 IID = Intrinsic::x86_avx512_pshuf_b_512;
3112 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3113 if (VecWidth == 128)
3114 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3115 else if (VecWidth == 256)
3116 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3117 else if (VecWidth == 512)
3118 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3121 }
else if (Name.starts_with(
"pmulh.w")) {
3122 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3123 "Unexpected intrinsic");
3126 }
else if (Name.starts_with(
"pmulhu.w")) {
3127 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3128 "Unexpected intrinsic");
3131 }
else if (Name.starts_with(
"pmaddw.d.")) {
3132 if (VecWidth == 128)
3133 IID = Intrinsic::x86_sse2_pmadd_wd;
3134 else if (VecWidth == 256)
3135 IID = Intrinsic::x86_avx2_pmadd_wd;
3136 else if (VecWidth == 512)
3137 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3140 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3141 if (VecWidth == 128)
3142 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3143 else if (VecWidth == 256)
3144 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3145 else if (VecWidth == 512)
3146 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3149 }
else if (Name.starts_with(
"packsswb.")) {
3150 if (VecWidth == 128)
3151 IID = Intrinsic::x86_sse2_packsswb_128;
3152 else if (VecWidth == 256)
3153 IID = Intrinsic::x86_avx2_packsswb;
3154 else if (VecWidth == 512)
3155 IID = Intrinsic::x86_avx512_packsswb_512;
3158 }
else if (Name.starts_with(
"packssdw.")) {
3159 if (VecWidth == 128)
3160 IID = Intrinsic::x86_sse2_packssdw_128;
3161 else if (VecWidth == 256)
3162 IID = Intrinsic::x86_avx2_packssdw;
3163 else if (VecWidth == 512)
3164 IID = Intrinsic::x86_avx512_packssdw_512;
3167 }
else if (Name.starts_with(
"packuswb.")) {
3168 if (VecWidth == 128)
3169 IID = Intrinsic::x86_sse2_packuswb_128;
3170 else if (VecWidth == 256)
3171 IID = Intrinsic::x86_avx2_packuswb;
3172 else if (VecWidth == 512)
3173 IID = Intrinsic::x86_avx512_packuswb_512;
3176 }
else if (Name.starts_with(
"packusdw.")) {
3177 if (VecWidth == 128)
3178 IID = Intrinsic::x86_sse41_packusdw;
3179 else if (VecWidth == 256)
3180 IID = Intrinsic::x86_avx2_packusdw;
3181 else if (VecWidth == 512)
3182 IID = Intrinsic::x86_avx512_packusdw_512;
3185 }
else if (Name.starts_with(
"vpermilvar.")) {
3186 if (VecWidth == 128 && EltWidth == 32)
3187 IID = Intrinsic::x86_avx_vpermilvar_ps;
3188 else if (VecWidth == 128 && EltWidth == 64)
3189 IID = Intrinsic::x86_avx_vpermilvar_pd;
3190 else if (VecWidth == 256 && EltWidth == 32)
3191 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3192 else if (VecWidth == 256 && EltWidth == 64)
3193 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3194 else if (VecWidth == 512 && EltWidth == 32)
3195 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3196 else if (VecWidth == 512 && EltWidth == 64)
3197 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3200 }
else if (Name ==
"cvtpd2dq.256") {
3201 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3202 }
else if (Name ==
"cvtpd2ps.256") {
3203 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3204 }
else if (Name ==
"cvttpd2dq.256") {
3205 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3206 }
else if (Name ==
"cvttps2dq.128") {
3207 IID = Intrinsic::x86_sse2_cvttps2dq;
3208 }
else if (Name ==
"cvttps2dq.256") {
3209 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3210 }
else if (Name.starts_with(
"permvar.")) {
3212 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3213 IID = Intrinsic::x86_avx2_permps;
3214 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3215 IID = Intrinsic::x86_avx2_permd;
3216 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3217 IID = Intrinsic::x86_avx512_permvar_df_256;
3218 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3219 IID = Intrinsic::x86_avx512_permvar_di_256;
3220 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3221 IID = Intrinsic::x86_avx512_permvar_sf_512;
3222 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3223 IID = Intrinsic::x86_avx512_permvar_si_512;
3224 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3225 IID = Intrinsic::x86_avx512_permvar_df_512;
3226 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3227 IID = Intrinsic::x86_avx512_permvar_di_512;
3228 else if (VecWidth == 128 && EltWidth == 16)
3229 IID = Intrinsic::x86_avx512_permvar_hi_128;
3230 else if (VecWidth == 256 && EltWidth == 16)
3231 IID = Intrinsic::x86_avx512_permvar_hi_256;
3232 else if (VecWidth == 512 && EltWidth == 16)
3233 IID = Intrinsic::x86_avx512_permvar_hi_512;
3234 else if (VecWidth == 128 && EltWidth == 8)
3235 IID = Intrinsic::x86_avx512_permvar_qi_128;
3236 else if (VecWidth == 256 && EltWidth == 8)
3237 IID = Intrinsic::x86_avx512_permvar_qi_256;
3238 else if (VecWidth == 512 && EltWidth == 8)
3239 IID = Intrinsic::x86_avx512_permvar_qi_512;
3242 }
else if (Name.starts_with(
"dbpsadbw.")) {
3243 if (VecWidth == 128)
3244 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3245 else if (VecWidth == 256)
3246 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3247 else if (VecWidth == 512)
3248 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3251 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3252 if (VecWidth == 128)
3253 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3254 else if (VecWidth == 256)
3255 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3256 else if (VecWidth == 512)
3257 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3260 }
else if (Name.starts_with(
"conflict.")) {
3261 if (Name[9] ==
'd' && VecWidth == 128)
3262 IID = Intrinsic::x86_avx512_conflict_d_128;
3263 else if (Name[9] ==
'd' && VecWidth == 256)
3264 IID = Intrinsic::x86_avx512_conflict_d_256;
3265 else if (Name[9] ==
'd' && VecWidth == 512)
3266 IID = Intrinsic::x86_avx512_conflict_d_512;
3267 else if (Name[9] ==
'q' && VecWidth == 128)
3268 IID = Intrinsic::x86_avx512_conflict_q_128;
3269 else if (Name[9] ==
'q' && VecWidth == 256)
3270 IID = Intrinsic::x86_avx512_conflict_q_256;
3271 else if (Name[9] ==
'q' && VecWidth == 512)
3272 IID = Intrinsic::x86_avx512_conflict_q_512;
3275 }
else if (Name.starts_with(
"pavg.")) {
3276 if (Name[5] ==
'b' && VecWidth == 128)
3277 IID = Intrinsic::x86_sse2_pavg_b;
3278 else if (Name[5] ==
'b' && VecWidth == 256)
3279 IID = Intrinsic::x86_avx2_pavg_b;
3280 else if (Name[5] ==
'b' && VecWidth == 512)
3281 IID = Intrinsic::x86_avx512_pavg_b_512;
3282 else if (Name[5] ==
'w' && VecWidth == 128)
3283 IID = Intrinsic::x86_sse2_pavg_w;
3284 else if (Name[5] ==
'w' && VecWidth == 256)
3285 IID = Intrinsic::x86_avx2_pavg_w;
3286 else if (Name[5] ==
'w' && VecWidth == 512)
3287 IID = Intrinsic::x86_avx512_pavg_w_512;
3296 Rep = Builder.CreateIntrinsic(IID, Args);
3307 if (AsmStr->find(
"mov\tfp") == 0 &&
3308 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3309 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3310 AsmStr->replace(Pos, 1,
";");
3318 assert(Result &&
"unsupported nvvm.{add,mul,fma}.* intrinsic");
3326 if (IsLegacyBF16Storage)
3327 Ty = Name.ends_with(
"bf16x2")
3329 :
static_cast<Type *
>(Builder.getBFloatTy());
3333 Args.push_back(IsLegacyBF16Storage ? Builder.CreateBitCast(Arg, Ty) : Arg);
3334 Args.push_back(Builder.getInt32(
static_cast<int>(
RoundingMode)));
3336 Value *Rep = Builder.CreateIntrinsic(Ty, IID, Args);
3337 return IsLegacyBF16Storage ? Builder.CreateBitCast(Rep, CI->
getType()) : Rep;
3342 Value *Rep =
nullptr;
3344 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3346 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3347 {Arg, Builder.getTrue()},
3349 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3350 Type *Ty = (Name ==
"abs.bf16")
3354 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3355 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3356 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3357 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3358 : Intrinsic::nvvm_fabs;
3359 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3360 }
else if (Name.consume_front(
"add.")) {
3363 }
else if (Name.consume_front(
"mul.")) {
3366 }
else if (Name.consume_front(
"ex2.approx.")) {
3368 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3369 : Intrinsic::nvvm_ex2_approx;
3370 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3371 }
else if (Name.consume_front(
"fma.")) {
3374 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3375 Name.starts_with(
"atomic.load.add.f64.p")) {
3378 Rep = Builder.CreateAtomicRMW(
3382 LLVMContext::MD_atomic_ignore_denormal_mode,
3387 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3388 Name.starts_with(
"atomic.load.dec.32.p")) {
3393 Rep = Builder.CreateAtomicRMW(
3397 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3403 Op.contains(
".cta.") ?
"block" :
"");
3404 if (
Op.starts_with(
"cas.")) {
3406 Value *Pair = Builder.CreateAtomicCmpXchg(
3409 Rep = Builder.CreateExtractValue(Pair, 0);
3427 "unexpected nvvm scoped atomic intrinsic");
3428 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3432 LLVMContext::MD_atomic_ignore_denormal_mode,
3436 }
else if (Name ==
"clz.ll") {
3439 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3440 {Arg, Builder.getFalse()},
3442 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3443 }
else if (Name ==
"popc.ll") {
3447 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3448 Arg,
nullptr,
"ctpop");
3449 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3450 }
else if (Name ==
"h2f") {
3452 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3453 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3454 }
else if (Name.consume_front(
"bitcast.") &&
3455 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3458 }
else if (Name ==
"rotate.b32") {
3461 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3462 {Arg, Arg, ShiftAmt});
3463 }
else if (Name ==
"rotate.b64") {
3467 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3468 {Arg, Arg, ZExtShiftAmt});
3469 }
else if (Name ==
"rotate.right.b64") {
3473 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3474 {Arg, Arg, ZExtShiftAmt});
3475 }
else if (Name ==
"swap.lo.hi.b64") {
3478 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3479 {Arg, Arg, Builder.getInt64(32)});
3480 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3483 Name.starts_with(
".to.gen"))) {
3485 }
else if (Name.consume_front(
"ldg.global")) {
3489 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3492 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3494 }
else if (Name ==
"tanh.approx.f32") {
3498 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3500 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3502 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3503 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3505 }
else if (Name ==
"barrier") {
3506 Rep = Builder.CreateIntrinsic(
3507 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3509 }
else if (Name ==
"barrier.sync") {
3510 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3512 }
else if (Name ==
"barrier.sync.cnt") {
3513 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3515 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3516 Name ==
"barrier0.or") {
3518 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3522 .
Case(
"barrier0.popc",
3523 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3524 .
Case(
"barrier0.and",
3525 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3526 .
Case(
"barrier0.or",
3527 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3528 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3529 Rep = Builder.CreateZExt(Bar, CI->
getType());
3543 ? Builder.CreateBitCast(Arg, NewType)
3546 Rep = Builder.CreateCall(NewFn, Args);
3547 if (
F->getReturnType()->isIntegerTy())
3548 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3558 Value *Rep =
nullptr;
3560 if (Name.starts_with(
"sse4a.movnt.")) {
3572 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3575 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3576 }
else if (Name.starts_with(
"avx.movnt.") ||
3577 Name.starts_with(
"avx512.storent.")) {
3589 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3590 }
else if (Name ==
"sse2.storel.dq") {
3595 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3596 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3597 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3598 }
else if (Name.starts_with(
"sse.storeu.") ||
3599 Name.starts_with(
"sse2.storeu.") ||
3600 Name.starts_with(
"avx.storeu.")) {
3603 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3604 }
else if (Name ==
"avx512.mask.store.ss") {
3608 }
else if (Name.starts_with(
"avx512.mask.store")) {
3610 bool Aligned = Name[17] !=
'u';
3613 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3616 bool CmpEq = Name[9] ==
'e';
3619 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3620 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3627 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3628 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3630 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3631 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3632 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3633 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3634 Name.starts_with(
"sse2.sqrt.p") ||
3635 Name.starts_with(
"sse.sqrt.p")) {
3636 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3637 {CI->getArgOperand(0)});
3638 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3642 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3643 : Intrinsic::x86_avx512_sqrt_pd_512;
3646 Rep = Builder.CreateIntrinsic(IID, Args);
3648 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3649 {CI->getArgOperand(0)});
3653 }
else if (Name.starts_with(
"avx512.ptestm") ||
3654 Name.starts_with(
"avx512.ptestnm")) {
3658 Rep = Builder.CreateAnd(Op0, Op1);
3664 Rep = Builder.CreateICmp(
Pred, Rep,
Zero);
3666 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3669 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3672 }
else if (Name.starts_with(
"avx512.kunpck")) {
3677 for (
unsigned i = 0; i != NumElts; ++i)
3686 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3687 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3688 }
else if (Name ==
"avx512.kand.w") {
3691 Rep = Builder.CreateAnd(
LHS,
RHS);
3692 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3693 }
else if (Name ==
"avx512.kandn.w") {
3696 LHS = Builder.CreateNot(
LHS);
3697 Rep = Builder.CreateAnd(
LHS,
RHS);
3698 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3699 }
else if (Name ==
"avx512.kor.w") {
3702 Rep = Builder.CreateOr(
LHS,
RHS);
3703 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3704 }
else if (Name ==
"avx512.kxor.w") {
3707 Rep = Builder.CreateXor(
LHS,
RHS);
3708 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3709 }
else if (Name ==
"avx512.kxnor.w") {
3712 LHS = Builder.CreateNot(
LHS);
3713 Rep = Builder.CreateXor(
LHS,
RHS);
3714 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3715 }
else if (Name ==
"avx512.knot.w") {
3717 Rep = Builder.CreateNot(Rep);
3718 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3719 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3722 Rep = Builder.CreateOr(
LHS,
RHS);
3723 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3725 if (Name[14] ==
'c')
3729 Rep = Builder.CreateICmpEQ(Rep,
C);
3730 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3731 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3732 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3733 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3734 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3737 ConstantInt::get(I32Ty, 0));
3739 ConstantInt::get(I32Ty, 0));
3741 if (Name.contains(
".add."))
3742 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3743 else if (Name.contains(
".sub."))
3744 EltOp = Builder.CreateFSub(Elt0, Elt1);
3745 else if (Name.contains(
".mul."))
3746 EltOp = Builder.CreateFMul(Elt0, Elt1);
3748 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3749 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3750 ConstantInt::get(I32Ty, 0));
3751 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3753 bool CmpEq = Name[16] ==
'e';
3755 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3757 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3764 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3767 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3770 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3777 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3779 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3780 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3782 if (VecWidth == 128 && EltWidth == 32)
3783 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3784 else if (VecWidth == 256 && EltWidth == 32)
3785 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3786 else if (VecWidth == 512 && EltWidth == 32)
3787 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3788 else if (VecWidth == 128 && EltWidth == 64)
3789 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3790 else if (VecWidth == 256 && EltWidth == 64)
3791 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3792 else if (VecWidth == 512 && EltWidth == 64)
3793 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3800 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3803 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3804 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3806 if (VecWidth == 128 && EltWidth == 32)
3807 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3808 else if (VecWidth == 256 && EltWidth == 32)
3809 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3810 else if (VecWidth == 512 && EltWidth == 32)
3811 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3812 else if (VecWidth == 128 && EltWidth == 64)
3813 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3814 else if (VecWidth == 256 && EltWidth == 64)
3815 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3816 else if (VecWidth == 512 && EltWidth == 64)
3817 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3822 if (VecWidth == 512)
3824 Args.push_back(Mask);
3826 Rep = Builder.CreateIntrinsic(IID, Args);
3827 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3831 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3834 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3835 Name.starts_with(
"avx512.cvtw2mask.") ||
3836 Name.starts_with(
"avx512.cvtd2mask.") ||
3837 Name.starts_with(
"avx512.cvtq2mask.")) {
3842 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3843 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3844 Name.starts_with(
"avx512.mask.pabs")) {
3846 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3847 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3848 Name.starts_with(
"avx512.mask.pmaxs")) {
3850 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3851 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3852 Name.starts_with(
"avx512.mask.pmaxu")) {
3854 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3855 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3856 Name.starts_with(
"avx512.mask.pmins")) {
3858 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3859 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3860 Name.starts_with(
"avx512.mask.pminu")) {
3862 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3863 Name.starts_with(
"avx512.pmulh.w")) {
3865 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3866 Name.starts_with(
"avx512.pmulhu.w")) {
3868 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3869 Name ==
"avx512.pmulu.dq.512" ||
3870 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3872 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3873 Name ==
"avx512.pmul.dq.512" ||
3874 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3876 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3877 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3882 }
else if (Name ==
"avx512.cvtusi2sd") {
3887 }
else if (Name ==
"sse2.cvtss2sd") {
3889 Rep = Builder.CreateFPExt(
3892 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3893 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3894 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3895 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3896 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3897 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3898 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3899 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3900 Name ==
"avx512.mask.cvtqq2ps.256" ||
3901 Name ==
"avx512.mask.cvtqq2ps.512" ||
3902 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3903 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3904 Name ==
"avx.cvt.ps2.pd.256" ||
3905 Name ==
"avx512.mask.cvtps2pd.128" ||
3906 Name ==
"avx512.mask.cvtps2pd.256") {
3911 unsigned NumDstElts = DstTy->getNumElements();
3912 if (NumDstElts < SrcTy->getNumElements()) {
3913 assert(NumDstElts == 2 &&
"Unexpected vector size");
3914 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3917 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3918 bool IsUnsigned = Name.contains(
"cvtu");
3920 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3924 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3925 : Intrinsic::x86_avx512_sitofp_round;
3926 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3929 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3930 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3936 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3937 Name.starts_with(
"vcvtph2ps.")) {
3941 unsigned NumDstElts = DstTy->getNumElements();
3942 if (NumDstElts != SrcTy->getNumElements()) {
3943 assert(NumDstElts == 4 &&
"Unexpected vector size");
3944 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3946 Rep = Builder.CreateBitCast(
3948 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3952 }
else if (Name.starts_with(
"avx512.mask.load")) {
3954 bool Aligned = Name[16] !=
'u';
3957 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3961 ResultTy->getNumElements());
3962 Rep = Builder.CreateIntrinsic(
3963 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3965 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3971 Rep = Builder.CreateIntrinsic(
3972 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3974 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3975 Name.starts_with(
"avx512.mask.expand.")) {
3979 ResultTy->getNumElements());
3981 bool IsCompress = Name[12] ==
'c';
3982 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3983 : Intrinsic::x86_avx512_mask_expand;
3984 Rep = Builder.CreateIntrinsic(
3986 }
else if (Name.starts_with(
"xop.vpcom")) {
3988 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3989 Name.ends_with(
"uq"))
3991 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3992 Name.ends_with(
"d") || Name.ends_with(
"q"))
4001 Name = Name.substr(9);
4002 if (Name.starts_with(
"lt"))
4004 else if (Name.starts_with(
"le"))
4006 else if (Name.starts_with(
"gt"))
4008 else if (Name.starts_with(
"ge"))
4010 else if (Name.starts_with(
"eq"))
4012 else if (Name.starts_with(
"ne"))
4014 else if (Name.starts_with(
"false"))
4016 else if (Name.starts_with(
"true"))
4023 }
else if (Name.starts_with(
"xop.vpcmov")) {
4025 Value *NotSel = Builder.CreateNot(Sel);
4028 Rep = Builder.CreateOr(Sel0, Sel1);
4029 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
4030 Name.starts_with(
"avx512.mask.prol")) {
4032 }
else if (Name.starts_with(
"avx512.pror") ||
4033 Name.starts_with(
"avx512.mask.pror")) {
4035 }
else if (Name.starts_with(
"avx512.vpshld.") ||
4036 Name.starts_with(
"avx512.mask.vpshld") ||
4037 Name.starts_with(
"avx512.maskz.vpshld")) {
4038 bool ZeroMask = Name[11] ==
'z';
4040 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
4041 Name.starts_with(
"avx512.mask.vpshrd") ||
4042 Name.starts_with(
"avx512.maskz.vpshrd")) {
4043 bool ZeroMask = Name[11] ==
'z';
4045 }
else if (Name ==
"sse42.crc32.64.8") {
4048 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
4050 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
4051 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
4052 Name.starts_with(
"avx512.vbroadcast.s")) {
4055 Type *EltTy = VecTy->getElementType();
4056 unsigned EltNum = VecTy->getNumElements();
4060 for (
unsigned I = 0;
I < EltNum; ++
I)
4061 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
4062 }
else if (Name.starts_with(
"sse41.pmovsx") ||
4063 Name.starts_with(
"sse41.pmovzx") ||
4064 Name.starts_with(
"avx2.pmovsx") ||
4065 Name.starts_with(
"avx2.pmovzx") ||
4066 Name.starts_with(
"avx512.mask.pmovsx") ||
4067 Name.starts_with(
"avx512.mask.pmovzx")) {
4069 unsigned NumDstElts = DstTy->getNumElements();
4073 for (
unsigned i = 0; i != NumDstElts; ++i)
4078 bool DoSext = Name.contains(
"pmovsx");
4080 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
4085 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
4086 Name ==
"avx512.mask.pmov.qd.512" ||
4087 Name ==
"avx512.mask.pmov.wb.256" ||
4088 Name ==
"avx512.mask.pmov.wb.512") {
4093 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4094 Name ==
"avx2.vbroadcasti128") {
4100 if (NumSrcElts == 2)
4103 Rep = Builder.CreateShuffleVector(
Load,
4105 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4106 Name.starts_with(
"avx512.mask.shuf.f")) {
4111 unsigned ControlBitsMask = NumLanes - 1;
4112 unsigned NumControlBits = NumLanes / 2;
4115 for (
unsigned l = 0; l != NumLanes; ++l) {
4116 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4118 if (l >= NumLanes / 2)
4119 LaneMask += NumLanes;
4120 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4121 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4127 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4128 Name.starts_with(
"avx512.mask.broadcasti")) {
4131 unsigned NumDstElts =
4135 for (
unsigned i = 0; i != NumDstElts; ++i)
4136 ShuffleMask[i] = i % NumSrcElts;
4142 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4143 Name.starts_with(
"avx2.vbroadcast") ||
4144 Name.starts_with(
"avx512.pbroadcast") ||
4145 Name.starts_with(
"avx512.mask.broadcast.s")) {
4152 Rep = Builder.CreateShuffleVector(
Op, M);
4157 }
else if (Name.starts_with(
"sse2.padds.") ||
4158 Name.starts_with(
"avx2.padds.") ||
4159 Name.starts_with(
"avx512.padds.") ||
4160 Name.starts_with(
"avx512.mask.padds.")) {
4162 }
else if (Name.starts_with(
"sse2.psubs.") ||
4163 Name.starts_with(
"avx2.psubs.") ||
4164 Name.starts_with(
"avx512.psubs.") ||
4165 Name.starts_with(
"avx512.mask.psubs.")) {
4167 }
else if (Name.starts_with(
"sse2.paddus.") ||
4168 Name.starts_with(
"avx2.paddus.") ||
4169 Name.starts_with(
"avx512.mask.paddus.")) {
4171 }
else if (Name.starts_with(
"sse2.psubus.") ||
4172 Name.starts_with(
"avx2.psubus.") ||
4173 Name.starts_with(
"avx512.mask.psubus.")) {
4175 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4180 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4184 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4189 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4194 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4195 Name ==
"avx512.psll.dq.512") {
4199 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4200 Name ==
"avx512.psrl.dq.512") {
4204 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4205 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4206 Name.starts_with(
"avx2.pblendd.")) {
4211 unsigned NumElts = VecTy->getNumElements();
4214 for (
unsigned i = 0; i != NumElts; ++i)
4215 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4217 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4218 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4219 Name ==
"avx2.vinserti128" ||
4220 Name.starts_with(
"avx512.mask.insert")) {
4224 unsigned DstNumElts =
4226 unsigned SrcNumElts =
4228 unsigned Scale = DstNumElts / SrcNumElts;
4235 for (
unsigned i = 0; i != SrcNumElts; ++i)
4237 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4238 Idxs[i] = SrcNumElts;
4239 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4253 for (
unsigned i = 0; i != DstNumElts; ++i)
4256 for (
unsigned i = 0; i != SrcNumElts; ++i)
4257 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4258 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4264 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4265 Name ==
"avx2.vextracti128" ||
4266 Name.starts_with(
"avx512.mask.vextract")) {
4269 unsigned DstNumElts =
4271 unsigned SrcNumElts =
4273 unsigned Scale = SrcNumElts / DstNumElts;
4280 for (
unsigned i = 0; i != DstNumElts; ++i) {
4281 Idxs[i] = i + (
Imm * DstNumElts);
4283 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4289 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4290 Name.starts_with(
"avx512.mask.perm.di.")) {
4294 unsigned NumElts = VecTy->getNumElements();
4297 for (
unsigned i = 0; i != NumElts; ++i)
4298 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4300 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4305 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4317 unsigned HalfSize = NumElts / 2;
4329 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4330 for (
unsigned i = 0; i < HalfSize; ++i)
4331 ShuffleMask[i] = StartIndex + i;
4334 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4335 for (
unsigned i = 0; i < HalfSize; ++i)
4336 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4338 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4340 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4341 Name.starts_with(
"avx512.mask.vpermil.p") ||
4342 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4346 unsigned NumElts = VecTy->getNumElements();
4348 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4349 unsigned IdxMask = ((1 << IdxSize) - 1);
4355 for (
unsigned i = 0; i != NumElts; ++i)
4356 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4358 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4363 }
else if (Name ==
"sse2.pshufl.w" ||
4364 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4369 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4373 for (
unsigned l = 0; l != NumElts; l += 8) {
4374 for (
unsigned i = 0; i != 4; ++i)
4375 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4376 for (
unsigned i = 4; i != 8; ++i)
4377 Idxs[i + l] = i + l;
4380 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4385 }
else if (Name ==
"sse2.pshufh.w" ||
4386 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4391 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4395 for (
unsigned l = 0; l != NumElts; l += 8) {
4396 for (
unsigned i = 0; i != 4; ++i)
4397 Idxs[i + l] = i + l;
4398 for (
unsigned i = 0; i != 4; ++i)
4399 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4402 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4407 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4414 unsigned HalfLaneElts = NumLaneElts / 2;
4417 for (
unsigned i = 0; i != NumElts; ++i) {
4419 Idxs[i] = i - (i % NumLaneElts);
4421 if ((i % NumLaneElts) >= HalfLaneElts)
4425 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4428 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4432 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4433 Name.starts_with(
"avx512.mask.movshdup") ||
4434 Name.starts_with(
"avx512.mask.movsldup")) {
4440 if (Name.starts_with(
"avx512.mask.movshdup."))
4444 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4445 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4446 Idxs[i + l + 0] = i + l +
Offset;
4447 Idxs[i + l + 1] = i + l +
Offset;
4450 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4454 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4455 Name.starts_with(
"avx512.mask.unpckl.")) {
4462 for (
int l = 0; l != NumElts; l += NumLaneElts)
4463 for (
int i = 0; i != NumLaneElts; ++i)
4464 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4466 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4470 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4471 Name.starts_with(
"avx512.mask.unpckh.")) {
4478 for (
int l = 0; l != NumElts; l += NumLaneElts)
4479 for (
int i = 0; i != NumLaneElts; ++i)
4480 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4482 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4486 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4487 Name.starts_with(
"avx512.mask.pand.")) {
4490 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4492 Rep = Builder.CreateBitCast(Rep, FTy);
4495 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4496 Name.starts_with(
"avx512.mask.pandn.")) {
4499 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4500 Rep = Builder.CreateAnd(Rep,
4502 Rep = Builder.CreateBitCast(Rep, FTy);
4505 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4506 Name.starts_with(
"avx512.mask.por.")) {
4509 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4511 Rep = Builder.CreateBitCast(Rep, FTy);
4514 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4515 Name.starts_with(
"avx512.mask.pxor.")) {
4518 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4520 Rep = Builder.CreateBitCast(Rep, FTy);
4523 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4527 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4531 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4535 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4536 if (Name.ends_with(
".512")) {
4538 if (Name[17] ==
's')
4539 IID = Intrinsic::x86_avx512_add_ps_512;
4541 IID = Intrinsic::x86_avx512_add_pd_512;
4543 Rep = Builder.CreateIntrinsic(
4551 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4552 if (Name.ends_with(
".512")) {
4554 if (Name[17] ==
's')
4555 IID = Intrinsic::x86_avx512_div_ps_512;
4557 IID = Intrinsic::x86_avx512_div_pd_512;
4559 Rep = Builder.CreateIntrinsic(
4567 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4568 if (Name.ends_with(
".512")) {
4570 if (Name[17] ==
's')
4571 IID = Intrinsic::x86_avx512_mul_ps_512;
4573 IID = Intrinsic::x86_avx512_mul_pd_512;
4575 Rep = Builder.CreateIntrinsic(
4583 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4584 if (Name.ends_with(
".512")) {
4586 if (Name[17] ==
's')
4587 IID = Intrinsic::x86_avx512_sub_ps_512;
4589 IID = Intrinsic::x86_avx512_sub_pd_512;
4591 Rep = Builder.CreateIntrinsic(
4599 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4600 Name.starts_with(
"avx512.mask.min.p")) &&
4601 Name.drop_front(18) ==
".512") {
4602 bool IsDouble = Name[17] ==
'd';
4603 bool IsMin = Name[13] ==
'i';
4605 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4606 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4609 Rep = Builder.CreateIntrinsic(
4614 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4616 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4617 {CI->getArgOperand(0), Builder.getInt1(false)});
4620 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4621 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4622 bool IsVariable = Name[16] ==
'v';
4623 char Size = Name[16] ==
'.' ? Name[17]
4624 : Name[17] ==
'.' ? Name[18]
4625 : Name[18] ==
'.' ? Name[19]
4629 if (IsVariable && Name[17] !=
'.') {
4630 if (
Size ==
'd' && Name[17] ==
'2')
4631 IID = Intrinsic::x86_avx2_psllv_q;
4632 else if (
Size ==
'd' && Name[17] ==
'4')
4633 IID = Intrinsic::x86_avx2_psllv_q_256;
4634 else if (
Size ==
's' && Name[17] ==
'4')
4635 IID = Intrinsic::x86_avx2_psllv_d;
4636 else if (
Size ==
's' && Name[17] ==
'8')
4637 IID = Intrinsic::x86_avx2_psllv_d_256;
4638 else if (
Size ==
'h' && Name[17] ==
'8')
4639 IID = Intrinsic::x86_avx512_psllv_w_128;
4640 else if (
Size ==
'h' && Name[17] ==
'1')
4641 IID = Intrinsic::x86_avx512_psllv_w_256;
4642 else if (Name[17] ==
'3' && Name[18] ==
'2')
4643 IID = Intrinsic::x86_avx512_psllv_w_512;
4646 }
else if (Name.ends_with(
".128")) {
4648 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4649 : Intrinsic::x86_sse2_psll_d;
4650 else if (
Size ==
'q')
4651 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4652 : Intrinsic::x86_sse2_psll_q;
4653 else if (
Size ==
'w')
4654 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4655 : Intrinsic::x86_sse2_psll_w;
4658 }
else if (Name.ends_with(
".256")) {
4660 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4661 : Intrinsic::x86_avx2_psll_d;
4662 else if (
Size ==
'q')
4663 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4664 : Intrinsic::x86_avx2_psll_q;
4665 else if (
Size ==
'w')
4666 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4667 : Intrinsic::x86_avx2_psll_w;
4672 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4673 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4674 : Intrinsic::x86_avx512_psll_d_512;
4675 else if (
Size ==
'q')
4676 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4677 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4678 : Intrinsic::x86_avx512_psll_q_512;
4679 else if (
Size ==
'w')
4680 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4681 : Intrinsic::x86_avx512_psll_w_512;
4687 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4688 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4689 bool IsVariable = Name[16] ==
'v';
4690 char Size = Name[16] ==
'.' ? Name[17]
4691 : Name[17] ==
'.' ? Name[18]
4692 : Name[18] ==
'.' ? Name[19]
4696 if (IsVariable && Name[17] !=
'.') {
4697 if (
Size ==
'd' && Name[17] ==
'2')
4698 IID = Intrinsic::x86_avx2_psrlv_q;
4699 else if (
Size ==
'd' && Name[17] ==
'4')
4700 IID = Intrinsic::x86_avx2_psrlv_q_256;
4701 else if (
Size ==
's' && Name[17] ==
'4')
4702 IID = Intrinsic::x86_avx2_psrlv_d;
4703 else if (
Size ==
's' && Name[17] ==
'8')
4704 IID = Intrinsic::x86_avx2_psrlv_d_256;
4705 else if (
Size ==
'h' && Name[17] ==
'8')
4706 IID = Intrinsic::x86_avx512_psrlv_w_128;
4707 else if (
Size ==
'h' && Name[17] ==
'1')
4708 IID = Intrinsic::x86_avx512_psrlv_w_256;
4709 else if (Name[17] ==
'3' && Name[18] ==
'2')
4710 IID = Intrinsic::x86_avx512_psrlv_w_512;
4713 }
else if (Name.ends_with(
".128")) {
4715 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4716 : Intrinsic::x86_sse2_psrl_d;
4717 else if (
Size ==
'q')
4718 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4719 : Intrinsic::x86_sse2_psrl_q;
4720 else if (
Size ==
'w')
4721 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4722 : Intrinsic::x86_sse2_psrl_w;
4725 }
else if (Name.ends_with(
".256")) {
4727 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4728 : Intrinsic::x86_avx2_psrl_d;
4729 else if (
Size ==
'q')
4730 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4731 : Intrinsic::x86_avx2_psrl_q;
4732 else if (
Size ==
'w')
4733 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4734 : Intrinsic::x86_avx2_psrl_w;
4739 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4740 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4741 : Intrinsic::x86_avx512_psrl_d_512;
4742 else if (
Size ==
'q')
4743 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4744 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4745 : Intrinsic::x86_avx512_psrl_q_512;
4746 else if (
Size ==
'w')
4747 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4748 : Intrinsic::x86_avx512_psrl_w_512;
4754 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4755 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4756 bool IsVariable = Name[16] ==
'v';
4757 char Size = Name[16] ==
'.' ? Name[17]
4758 : Name[17] ==
'.' ? Name[18]
4759 : Name[18] ==
'.' ? Name[19]
4763 if (IsVariable && Name[17] !=
'.') {
4764 if (
Size ==
's' && Name[17] ==
'4')
4765 IID = Intrinsic::x86_avx2_psrav_d;
4766 else if (
Size ==
's' && Name[17] ==
'8')
4767 IID = Intrinsic::x86_avx2_psrav_d_256;
4768 else if (
Size ==
'h' && Name[17] ==
'8')
4769 IID = Intrinsic::x86_avx512_psrav_w_128;
4770 else if (
Size ==
'h' && Name[17] ==
'1')
4771 IID = Intrinsic::x86_avx512_psrav_w_256;
4772 else if (Name[17] ==
'3' && Name[18] ==
'2')
4773 IID = Intrinsic::x86_avx512_psrav_w_512;
4776 }
else if (Name.ends_with(
".128")) {
4778 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4779 : Intrinsic::x86_sse2_psra_d;
4780 else if (
Size ==
'q')
4781 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4782 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4783 : Intrinsic::x86_avx512_psra_q_128;
4784 else if (
Size ==
'w')
4785 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4786 : Intrinsic::x86_sse2_psra_w;
4789 }
else if (Name.ends_with(
".256")) {
4791 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4792 : Intrinsic::x86_avx2_psra_d;
4793 else if (
Size ==
'q')
4794 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4795 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4796 : Intrinsic::x86_avx512_psra_q_256;
4797 else if (
Size ==
'w')
4798 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4799 : Intrinsic::x86_avx2_psra_w;
4804 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4805 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4806 : Intrinsic::x86_avx512_psra_d_512;
4807 else if (
Size ==
'q')
4808 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4809 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4810 : Intrinsic::x86_avx512_psra_q_512;
4811 else if (
Size ==
'w')
4812 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4813 : Intrinsic::x86_avx512_psra_w_512;
4819 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4821 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4823 }
else if (Name.ends_with(
".movntdqa")) {
4827 LoadInst *LI = Builder.CreateAlignedLoad(
4832 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4833 Name.starts_with(
"fma.vfmsub.") ||
4834 Name.starts_with(
"fma.vfnmadd.") ||
4835 Name.starts_with(
"fma.vfnmsub.")) {
4836 bool NegMul = Name[6] ==
'n';
4837 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4838 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4849 if (NegMul && !IsScalar)
4850 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4851 if (NegMul && IsScalar)
4852 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4854 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4856 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4860 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4868 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4872 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4873 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4874 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4875 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4876 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4877 bool IsMask3 = Name[11] ==
'3';
4878 bool IsMaskZ = Name[11] ==
'z';
4880 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4881 bool NegMul = Name[2] ==
'n';
4882 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4888 if (NegMul && (IsMask3 || IsMaskZ))
4889 A = Builder.CreateFNeg(
A);
4890 if (NegMul && !(IsMask3 || IsMaskZ))
4891 B = Builder.CreateFNeg(
B);
4893 C = Builder.CreateFNeg(
C);
4895 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4896 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4897 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4904 if (Name.back() ==
'd')
4905 IID = Intrinsic::x86_avx512_vfmadd_f64;
4907 IID = Intrinsic::x86_avx512_vfmadd_f32;
4908 Rep = Builder.CreateIntrinsic(IID,
Ops);
4910 Rep = Builder.CreateFMA(
A,
B,
C);
4919 if (NegAcc && IsMask3)
4924 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4926 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4927 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4928 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4929 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4930 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4931 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4932 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4933 bool IsMask3 = Name[11] ==
'3';
4934 bool IsMaskZ = Name[11] ==
'z';
4936 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4937 bool NegMul = Name[2] ==
'n';
4938 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4944 if (NegMul && (IsMask3 || IsMaskZ))
4945 A = Builder.CreateFNeg(
A);
4946 if (NegMul && !(IsMask3 || IsMaskZ))
4947 B = Builder.CreateFNeg(
B);
4949 C = Builder.CreateFNeg(
C);
4956 if (Name[Name.size() - 5] ==
's')
4957 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4959 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4963 Rep = Builder.CreateFMA(
A,
B,
C);
4971 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4975 if (VecWidth == 128 && EltWidth == 32)
4976 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4977 else if (VecWidth == 256 && EltWidth == 32)
4978 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4979 else if (VecWidth == 128 && EltWidth == 64)
4980 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4981 else if (VecWidth == 256 && EltWidth == 64)
4982 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4988 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4989 Rep = Builder.CreateIntrinsic(IID,
Ops);
4990 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4991 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4992 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4993 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4994 bool IsMask3 = Name[11] ==
'3';
4995 bool IsMaskZ = Name[11] ==
'z';
4997 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4998 bool IsSubAdd = Name[3] ==
's';
5002 if (Name[Name.size() - 5] ==
's')
5003 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
5005 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
5010 Ops[2] = Builder.CreateFNeg(
Ops[2]);
5012 Rep = Builder.CreateIntrinsic(IID,
Ops);
5021 Value *Odd = Builder.CreateCall(FMA,
Ops);
5022 Ops[2] = Builder.CreateFNeg(
Ops[2]);
5023 Value *Even = Builder.CreateCall(FMA,
Ops);
5029 for (
int i = 0; i != NumElts; ++i)
5030 Idxs[i] = i + (i % 2) * NumElts;
5032 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
5040 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
5041 Name.starts_with(
"avx512.maskz.pternlog.")) {
5042 bool ZeroMask = Name[11] ==
'z';
5046 if (VecWidth == 128 && EltWidth == 32)
5047 IID = Intrinsic::x86_avx512_pternlog_d_128;
5048 else if (VecWidth == 256 && EltWidth == 32)
5049 IID = Intrinsic::x86_avx512_pternlog_d_256;
5050 else if (VecWidth == 512 && EltWidth == 32)
5051 IID = Intrinsic::x86_avx512_pternlog_d_512;
5052 else if (VecWidth == 128 && EltWidth == 64)
5053 IID = Intrinsic::x86_avx512_pternlog_q_128;
5054 else if (VecWidth == 256 && EltWidth == 64)
5055 IID = Intrinsic::x86_avx512_pternlog_q_256;
5056 else if (VecWidth == 512 && EltWidth == 64)
5057 IID = Intrinsic::x86_avx512_pternlog_q_512;
5063 Rep = Builder.CreateIntrinsic(IID, Args);
5067 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
5068 Name.starts_with(
"avx512.maskz.vpmadd52")) {
5069 bool ZeroMask = Name[11] ==
'z';
5070 bool High = Name[20] ==
'h' || Name[21] ==
'h';
5073 if (VecWidth == 128 && !
High)
5074 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
5075 else if (VecWidth == 256 && !
High)
5076 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
5077 else if (VecWidth == 512 && !
High)
5078 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
5079 else if (VecWidth == 128 &&
High)
5080 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
5081 else if (VecWidth == 256 &&
High)
5082 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
5083 else if (VecWidth == 512 &&
High)
5084 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5090 Rep = Builder.CreateIntrinsic(IID, Args);
5094 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5095 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5096 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5097 bool ZeroMask = Name[11] ==
'z';
5098 bool IndexForm = Name[17] ==
'i';
5100 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5101 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5102 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5103 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5104 bool ZeroMask = Name[11] ==
'z';
5105 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5108 if (VecWidth == 128 && !IsSaturating)
5109 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5110 else if (VecWidth == 256 && !IsSaturating)
5111 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5112 else if (VecWidth == 512 && !IsSaturating)
5113 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5114 else if (VecWidth == 128 && IsSaturating)
5115 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5116 else if (VecWidth == 256 && IsSaturating)
5117 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5118 else if (VecWidth == 512 && IsSaturating)
5119 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5129 if (Args[1]->
getType()->isVectorTy() &&
5132 ->isIntegerTy(32) &&
5133 Args[2]->
getType()->isVectorTy() &&
5136 ->isIntegerTy(32)) {
5137 Type *NewArgType =
nullptr;
5138 if (VecWidth == 128)
5140 else if (VecWidth == 256)
5142 else if (VecWidth == 512)
5148 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5149 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5152 Rep = Builder.CreateIntrinsic(IID, Args);
5156 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5157 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5158 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5159 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5160 bool ZeroMask = Name[11] ==
'z';
5161 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5164 if (VecWidth == 128 && !IsSaturating)
5165 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5166 else if (VecWidth == 256 && !IsSaturating)
5167 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5168 else if (VecWidth == 512 && !IsSaturating)
5169 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5170 else if (VecWidth == 128 && IsSaturating)
5171 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5172 else if (VecWidth == 256 && IsSaturating)
5173 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5174 else if (VecWidth == 512 && IsSaturating)
5175 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5185 if (Args[1]->
getType()->isVectorTy() &&
5188 ->isIntegerTy(32) &&
5189 Args[2]->
getType()->isVectorTy() &&
5192 ->isIntegerTy(32)) {
5193 Type *NewArgType =
nullptr;
5194 if (VecWidth == 128)
5196 else if (VecWidth == 256)
5198 else if (VecWidth == 512)
5204 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5205 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5208 Rep = Builder.CreateIntrinsic(IID, Args);
5212 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5213 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5214 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5216 if (Name[0] ==
'a' && Name.back() ==
'2')
5217 IID = Intrinsic::x86_addcarry_32;
5218 else if (Name[0] ==
'a' && Name.back() ==
'4')
5219 IID = Intrinsic::x86_addcarry_64;
5220 else if (Name[0] ==
's' && Name.back() ==
'2')
5221 IID = Intrinsic::x86_subborrow_32;
5222 else if (Name[0] ==
's' && Name.back() ==
'4')
5223 IID = Intrinsic::x86_subborrow_64;
5230 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5233 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5236 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5240 }
else if (Name.starts_with(
"avx512.mask.") &&
5243 }
else if (Name.starts_with(
"bmi.pdep.")) {
5245 }
else if (Name.starts_with(
"bmi.pext.")) {
5255 if (Name.starts_with(
"neon.bfcvt")) {
5256 if (Name.starts_with(
"neon.bfcvtn2")) {
5258 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5260 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5261 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5264 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5265 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5267 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5271 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5272 return Builder.CreateShuffleVector(
5275 return Builder.CreateFPTrunc(CI->
getOperand(0),
5278 }
else if (Name.starts_with(
"sve.fcvt")) {
5281 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5282 .
Case(
"sve.fcvtnt.bf16f32",
5283 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5295 if (Args[1]->
getType() != BadPredTy)
5298 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5299 BadPredTy, Args[1]);
5300 Args[1] = Builder.CreateIntrinsic(
5301 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5303 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5307 if (Name ==
"neon.vcvtfp2hf")
5308 return Builder.CreateBitCast(
5309 Builder.CreateFPTrunc(
5313 if (Name ==
"neon.vcvthf2fp")
5314 return Builder.CreateFPExt(
5315 Builder.CreateBitCast(
5325 if (Name ==
"mve.vctp64.old") {
5328 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5331 Value *
C1 = Builder.CreateIntrinsic(
5332 Intrinsic::arm_mve_pred_v2i,
5334 return Builder.CreateIntrinsic(
5335 Intrinsic::arm_mve_pred_i2v,
5337 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5338 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5339 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5340 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5342 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5343 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5344 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5345 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5347 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5348 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5349 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5350 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5351 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5352 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5353 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5354 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5355 std::vector<Type *> Tys;
5359 case Intrinsic::arm_mve_mull_int_predicated:
5360 case Intrinsic::arm_mve_vqdmull_predicated:
5361 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5364 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5365 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5366 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5370 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5374 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5378 case Intrinsic::arm_cde_vcx1q_predicated:
5379 case Intrinsic::arm_cde_vcx1qa_predicated:
5380 case Intrinsic::arm_cde_vcx2q_predicated:
5381 case Intrinsic::arm_cde_vcx2qa_predicated:
5382 case Intrinsic::arm_cde_vcx3q_predicated:
5383 case Intrinsic::arm_cde_vcx3qa_predicated:
5390 std::vector<Value *>
Ops;
5392 Type *Ty =
Op->getType();
5393 if (Ty->getScalarSizeInBits() == 1) {
5394 Value *
C1 = Builder.CreateIntrinsic(
5395 Intrinsic::arm_mve_pred_v2i,
5397 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty},
C1);
5402 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5417 auto UpgradeLegacyWMMAIUIntrinsicCall =
5422 Args.push_back(Builder.getFalse());
5426 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5433 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5437 NewCall->copyMetadata(*CI);
5441 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5442 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5443 "intrinsic should have 7 arguments");
5446 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5448 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5449 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5450 "intrinsic should have 8 arguments");
5455 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5458 switch (
F->getIntrinsicID()) {
5461 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5462 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5463 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5464 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5465 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5466 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5481 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5484 F->getParent(),
F->getIntrinsicID(), Overloads);
5489 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5493 NewCall->copyMetadata(*CI);
5494 NewCall->takeName(CI);
5499 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5505 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5506 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5514 if (Name.starts_with(
"addrspacecast.nonnull")) {
5517 Value *ASC = Builder.CreateAddrSpaceCast(
5540 if (NumOperands < 3)
5553 bool IsVolatile =
false;
5557 if (NumOperands > 3)
5562 if (NumOperands > 5) {
5564 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5578 if (VT->getElementType()->isIntegerTy(16)) {
5581 Val = Builder.CreateBitCast(Val, AsBF16);
5589 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5591 unsigned AddrSpace = PtrTy->getAddressSpace();
5594 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5596 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5601 MDNode *RangeNotPrivate =
5604 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5610 return Builder.CreateBitCast(RMW, RetTy);
5631 return MAV->getMetadata();
5640 if (Name ==
"label") {
5642 }
else if (Name ==
"assign") {
5649 }
else if (Name ==
"declare") {
5653 }
else if (Name ==
"addr") {
5663 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5664 }
else if (Name ==
"value") {
5667 unsigned ExprOp = 2;
5682 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5690 int64_t OffsetVal =
Offset->getSExtValue();
5691 return Builder.CreateIntrinsic(OffsetVal >= 0
5692 ? Intrinsic::vector_splice_left
5693 : Intrinsic::vector_splice_right,
5695 {CI->getArgOperand(0), CI->getArgOperand(1),
5696 Builder.getInt32(std::abs(OffsetVal))});
5701 if (Name.starts_with(
"to.fp16")) {
5703 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5704 return Builder.CreateBitCast(Cast, CI->
getType());
5707 if (Name.starts_with(
"from.fp16")) {
5709 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5710 return Builder.CreateFPExt(Cast, CI->
getType());
5769 else if (Opcode == Instruction::ICmp)
5772 else if (Opcode == Instruction::FCmp)
5775 else if (Opcode == Instruction::Select)
5780 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5792 if (Defaults.empty())
5795 unsigned OldArgCount = CI->
arg_size();
5796 unsigned NewArgCount = NewFn->
arg_size();
5798 if (OldArgCount < FirstDefault)
5802 if (OldArgCount > NewArgCount)
5807 if (OldArgCount == NewArgCount) {
5819 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5820 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5821 "missing argument outside the default range");
5822 Type *ParamTy = NewFT->getParamType(Idx);
5827 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5833 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5864 if (!Name.consume_front(
"llvm."))
5867 bool IsX86 = Name.consume_front(
"x86.");
5868 bool IsNVVM = Name.consume_front(
"nvvm.");
5869 bool IsAArch64 = Name.consume_front(
"aarch64.");
5870 bool IsARM = Name.consume_front(
"arm.");
5871 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5872 bool IsDbg = Name.consume_front(
"dbg.");
5874 (Name.consume_front(
"experimental.vector.splice") ||
5875 Name.consume_front(
"vector.splice")) &&
5876 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5877 Value *Rep =
nullptr;
5879 if (!IsX86 && Name ==
"stackprotectorcheck") {
5881 }
else if (IsNVVM) {
5885 }
else if (IsAArch64) {
5889 }
else if (IsAMDGCN) {
5893 }
else if (IsOldSplice) {
5895 }
else if (Name.consume_front(
"convert.")) {
5897 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5912 const auto &DefaultCase = [&]() ->
void {
5920 "Unknown function for CallBase upgrade and isn't just a name change");
5928 "Return type must have changed");
5929 assert(OldST->getNumElements() ==
5931 "Must have same number of elements");
5934 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5937 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5938 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5939 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5960 case Intrinsic::arm_neon_vst1:
5961 case Intrinsic::arm_neon_vst2:
5962 case Intrinsic::arm_neon_vst3:
5963 case Intrinsic::arm_neon_vst4:
5964 case Intrinsic::arm_neon_vst2lane:
5965 case Intrinsic::arm_neon_vst3lane:
5966 case Intrinsic::arm_neon_vst4lane: {
5968 NewCall = Builder.CreateCall(NewFn, Args);
5971 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5972 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5973 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5978 NewCall = Builder.CreateCall(NewFn, Args);
5981 case Intrinsic::aarch64_sve_ld3_sret:
5982 case Intrinsic::aarch64_sve_ld4_sret:
5983 case Intrinsic::aarch64_sve_ld2_sret: {
5991 Name = Name.substr(5);
5998 unsigned MinElts = RetTy->getMinNumElements() /
N;
6000 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
6002 for (
unsigned I = 0;
I <
N;
I++) {
6003 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
6004 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
6010 case Intrinsic::coro_end_async:
6011 case Intrinsic::coro_end: {
6013 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
6015 NewCall = Builder.CreateCall(NewFn, Args);
6020 CI->
getModule(), Intrinsic::coro_is_in_ramp);
6021 Value *InRamp = Builder.CreateCall(IsInRamp);
6031 case Intrinsic::vector_extract: {
6033 Name = Name.substr(5);
6034 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
6039 unsigned MinElts = RetTy->getMinNumElements();
6042 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
6046 case Intrinsic::vector_insert: {
6048 Name = Name.substr(5);
6049 if (!Name.starts_with(
"aarch64.sve.tuple")) {
6053 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
6058 NewCall = Builder.CreateCall(
6062 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
6068 assert(
N > 1 &&
"Create is expected to be between 2-4");
6071 unsigned MinElts = RetTy->getMinNumElements() /
N;
6072 for (
unsigned I = 0;
I <
N;
I++) {
6074 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
6081 case Intrinsic::arm_neon_bfdot:
6082 case Intrinsic::arm_neon_bfmmla:
6083 case Intrinsic::arm_neon_bfmlalb:
6084 case Intrinsic::arm_neon_bfmlalt:
6085 case Intrinsic::aarch64_neon_bfdot:
6086 case Intrinsic::aarch64_neon_bfmmla:
6087 case Intrinsic::aarch64_neon_bfmlalb:
6088 case Intrinsic::aarch64_neon_bfmlalt: {
6091 "Mismatch between function args and call args");
6092 size_t OperandWidth =
6094 assert((OperandWidth == 64 || OperandWidth == 128) &&
6095 "Unexpected operand width");
6097 auto Iter = CI->
args().begin();
6098 Args.push_back(*Iter++);
6099 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6100 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6101 NewCall = Builder.CreateCall(NewFn, Args);
6105 case Intrinsic::bitreverse:
6106 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6109 case Intrinsic::ctlz:
6110 case Intrinsic::cttz: {
6117 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6121 case Intrinsic::objectsize: {
6122 Value *NullIsUnknownSize =
6126 NewCall = Builder.CreateCall(
6131 case Intrinsic::ctpop:
6132 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6134 case Intrinsic::dbg_value: {
6136 Name = Name.substr(5);
6138 if (Name.starts_with(
"dbg.addr")) {
6152 if (
Offset->isNullValue()) {
6153 NewCall = Builder.CreateCall(
6162 case Intrinsic::ptr_annotation:
6170 NewCall = Builder.CreateCall(
6179 case Intrinsic::var_annotation:
6186 NewCall = Builder.CreateCall(
6195 case Intrinsic::riscv_aes32dsi:
6196 case Intrinsic::riscv_aes32dsmi:
6197 case Intrinsic::riscv_aes32esi:
6198 case Intrinsic::riscv_aes32esmi:
6199 case Intrinsic::riscv_sm4ks:
6200 case Intrinsic::riscv_sm4ed: {
6210 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6211 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6217 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6218 Value *Res = NewCall;
6220 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6226 case Intrinsic::nvvm_mapa_shared_cluster: {
6230 Value *Res = NewCall;
6231 Res = Builder.CreateAddrSpaceCast(
6238 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6240 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6242 Args[0] = Builder.CreateAddrSpaceCast(
6246 Args.push_back(Builder.getInt32(0));
6248 NewCall = Builder.CreateCall(NewFn, Args);
6254 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6259 for (
unsigned I = 0;
I < 4; ++
I)
6261 Args.push_back(Builder.getInt32(0));
6262 Args.push_back(Builder.getInt32(0));
6265 Args.push_back(Builder.getInt1(
false));
6266 Args.push_back(Builder.getInt32(0));
6268 NewCall = Builder.CreateCall(NewFn, Args);
6274 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6277 Args[0] = Builder.CreateAddrSpaceCast(
6280 NewCall = Builder.CreateCall(NewFn, Args);
6287#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6288 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6290#undef G2S_CLUSTER_CASE
6295 Args[0] = Builder.CreateAddrSpaceCast(
6301 Args.push_back(Builder.getInt32(0));
6303 NewCall = Builder.CreateCall(NewFn, Args);
6310#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6311 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6319 "expected only the trailing flag_valid_pattern to be missing");
6320 Args.push_back(Builder.getInt32(0));
6322 NewCall = Builder.CreateCall(NewFn, Args);
6328#undef NVVM_TMA_G2S_MODES
6331 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6332 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6333 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6334 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6335 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6336 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6337 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6338 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6340 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6344 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6345 NewCall = Builder.CreateCall(NewFn, Args);
6349 case Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cluster:
6350 case Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cluster:
6351 case Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cluster:
6352 case Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cluster:
6353 case Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cluster:
6354 case Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cluster:
6355 case Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cluster:
6356 case Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cluster:
6357 case Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cluster:
6358 case Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cluster:
6359 case Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cluster:
6360 case Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cluster:
6361 case Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cluster:
6362 case Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cluster:
6363 case Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cluster:
6364 case Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cluster:
6365 case Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cluster:
6366 case Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cluster:
6367 case Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cluster:
6368 case Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cluster: {
6371 Args.push_back(Builder.getInt32(0));
6372 Args.push_back(Builder.getFalse());
6373 NewCall = Builder.CreateCall(NewFn, Args);
6379 case Intrinsic::nvvm_mbarrier_init: {
6383 if (Args.size() == 2)
6384 Args.push_back(Builder.getInt32(0));
6385 NewCall = Builder.CreateCall(NewFn, Args);
6388 case Intrinsic::riscv_sha256sig0:
6389 case Intrinsic::riscv_sha256sig1:
6390 case Intrinsic::riscv_sha256sum0:
6391 case Intrinsic::riscv_sha256sum1:
6392 case Intrinsic::riscv_sm3p0:
6393 case Intrinsic::riscv_sm3p1: {
6400 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6402 NewCall = Builder.CreateCall(NewFn, Arg);
6404 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6411 case Intrinsic::x86_xop_vfrcz_ss:
6412 case Intrinsic::x86_xop_vfrcz_sd:
6413 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6416 case Intrinsic::x86_xop_vpermil2pd:
6417 case Intrinsic::x86_xop_vpermil2ps:
6418 case Intrinsic::x86_xop_vpermil2pd_256:
6419 case Intrinsic::x86_xop_vpermil2ps_256: {
6423 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6424 NewCall = Builder.CreateCall(NewFn, Args);
6428 case Intrinsic::x86_sse41_ptestc:
6429 case Intrinsic::x86_sse41_ptestz:
6430 case Intrinsic::x86_sse41_ptestnzc: {
6444 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6445 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6447 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6451 case Intrinsic::x86_rdtscp: {
6457 NewCall = Builder.CreateCall(NewFn);
6459 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6462 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6470 case Intrinsic::x86_sse41_insertps:
6471 case Intrinsic::x86_sse41_dppd:
6472 case Intrinsic::x86_sse41_dpps:
6473 case Intrinsic::x86_sse41_mpsadbw:
6474 case Intrinsic::x86_avx_dp_ps_256:
6475 case Intrinsic::x86_avx2_mpsadbw: {
6481 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6482 NewCall = Builder.CreateCall(NewFn, Args);
6486 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6487 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6488 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6489 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6490 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6491 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6497 NewCall = Builder.CreateCall(NewFn, Args);
6506 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6507 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6508 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6509 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6510 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6511 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6515 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6516 Args[1] = Builder.CreateBitCast(
6519 NewCall = Builder.CreateCall(NewFn, Args);
6520 Value *Res = Builder.CreateBitCast(
6528 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6529 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6530 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6534 Args[1] = Builder.CreateBitCast(
6536 Args[2] = Builder.CreateBitCast(
6539 NewCall = Builder.CreateCall(NewFn, Args);
6543 case Intrinsic::amdgcn_asyncmark:
6544 case Intrinsic::amdgcn_wait_asyncmark: {
6548 Args.push_back(Builder.getInt32(0));
6553 NewCall = Builder.CreateCall(NewFn, Args, Bundles);
6561 case Intrinsic::thread_pointer: {
6562 NewCall = Builder.CreateCall(NewFn, {});
6566 case Intrinsic::memcpy:
6567 case Intrinsic::memmove:
6568 case Intrinsic::memset: {
6584 NewCall = Builder.CreateCall(NewFn, Args);
6587 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6588 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6589 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6594 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6597 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6601 case Intrinsic::masked_load:
6602 case Intrinsic::masked_gather:
6603 case Intrinsic::masked_store:
6604 case Intrinsic::masked_scatter: {
6610 auto GetMaybeAlign = [](
Value *
Op) {
6612 uint64_t Val = CI->getZExtValue();
6620 auto GetAlign = [&](
Value *
Op) {
6629 case Intrinsic::masked_load:
6630 NewCall = Builder.CreateMaskedLoad(
6634 case Intrinsic::masked_gather:
6635 NewCall = Builder.CreateMaskedGather(
6641 case Intrinsic::masked_store:
6642 NewCall = Builder.CreateMaskedStore(
6646 case Intrinsic::masked_scatter:
6647 NewCall = Builder.CreateMaskedScatter(
6649 DL.getValueOrABITypeAlignment(
6663 case Intrinsic::lifetime_start:
6664 case Intrinsic::lifetime_end: {
6676 NewCall = Builder.CreateLifetimeStart(Ptr);
6678 NewCall = Builder.CreateLifetimeEnd(Ptr);
6687 case Intrinsic::x86_avx512_vpdpbusd_128:
6688 case Intrinsic::x86_avx512_vpdpbusd_256:
6689 case Intrinsic::x86_avx512_vpdpbusd_512:
6690 case Intrinsic::x86_avx512_vpdpbusds_128:
6691 case Intrinsic::x86_avx512_vpdpbusds_256:
6692 case Intrinsic::x86_avx512_vpdpbusds_512:
6693 case Intrinsic::x86_avx2_vpdpbssd_128:
6694 case Intrinsic::x86_avx2_vpdpbssd_256:
6695 case Intrinsic::x86_avx10_vpdpbssd_512:
6696 case Intrinsic::x86_avx2_vpdpbssds_128:
6697 case Intrinsic::x86_avx2_vpdpbssds_256:
6698 case Intrinsic::x86_avx10_vpdpbssds_512:
6699 case Intrinsic::x86_avx2_vpdpbsud_128:
6700 case Intrinsic::x86_avx2_vpdpbsud_256:
6701 case Intrinsic::x86_avx10_vpdpbsud_512:
6702 case Intrinsic::x86_avx2_vpdpbsuds_128:
6703 case Intrinsic::x86_avx2_vpdpbsuds_256:
6704 case Intrinsic::x86_avx10_vpdpbsuds_512:
6705 case Intrinsic::x86_avx2_vpdpbuud_128:
6706 case Intrinsic::x86_avx2_vpdpbuud_256:
6707 case Intrinsic::x86_avx10_vpdpbuud_512:
6708 case Intrinsic::x86_avx2_vpdpbuuds_128:
6709 case Intrinsic::x86_avx2_vpdpbuuds_256:
6710 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6715 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6716 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6718 NewCall = Builder.CreateCall(NewFn, Args);
6721 case Intrinsic::x86_avx512_vpdpwssd_128:
6722 case Intrinsic::x86_avx512_vpdpwssd_256:
6723 case Intrinsic::x86_avx512_vpdpwssd_512:
6724 case Intrinsic::x86_avx512_vpdpwssds_128:
6725 case Intrinsic::x86_avx512_vpdpwssds_256:
6726 case Intrinsic::x86_avx512_vpdpwssds_512:
6727 case Intrinsic::x86_avx2_vpdpwsud_128:
6728 case Intrinsic::x86_avx2_vpdpwsud_256:
6729 case Intrinsic::x86_avx10_vpdpwsud_512:
6730 case Intrinsic::x86_avx2_vpdpwsuds_128:
6731 case Intrinsic::x86_avx2_vpdpwsuds_256:
6732 case Intrinsic::x86_avx10_vpdpwsuds_512:
6733 case Intrinsic::x86_avx2_vpdpwusd_128:
6734 case Intrinsic::x86_avx2_vpdpwusd_256:
6735 case Intrinsic::x86_avx10_vpdpwusd_512:
6736 case Intrinsic::x86_avx2_vpdpwusds_128:
6737 case Intrinsic::x86_avx2_vpdpwusds_256:
6738 case Intrinsic::x86_avx10_vpdpwusds_512:
6739 case Intrinsic::x86_avx2_vpdpwuud_128:
6740 case Intrinsic::x86_avx2_vpdpwuud_256:
6741 case Intrinsic::x86_avx10_vpdpwuud_512:
6742 case Intrinsic::x86_avx2_vpdpwuuds_128:
6743 case Intrinsic::x86_avx2_vpdpwuuds_256:
6744 case Intrinsic::x86_avx10_vpdpwuuds_512:
6749 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6750 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6752 NewCall = Builder.CreateCall(NewFn, Args);
6755 assert(NewCall &&
"Should have either set this variable or returned through "
6756 "the default case");
6763 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6777 F->eraseFromParent();
6783 if (NumOperands == 0)
6791 if (NumOperands == 3) {
6795 Metadata *Elts2[] = {ScalarType, ScalarType,
6811 if (NumOperands == 0 || NumOperands % 3 != 0)
6816 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6821 if (Upgraded ==
Tag)
6831 if (
Opc != Instruction::BitCast)
6835 Type *SrcTy = V->getType();
6852 if (
Opc != Instruction::BitCast)
6855 Type *SrcTy =
C->getType();
6872 if (Flag.getNumOperands() < 3)
6873 return std::nullopt;
6875 return Name->getString();
6876 return std::nullopt;
6882 if (!
getIROptions(M.getContext()).auto_upgrade_debug_info)
6890 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6891 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6892 if (auto Name = getModuleFlagNameSafely(*Flag))
6893 return *Name ==
"Debug Info Version";
6896 if (OpIt != ModFlags->op_end()) {
6897 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6904 bool BrokenDebugInfo =
false;
6907 if (!BrokenDebugInfo)
6913 M.getContext().diagnose(Diag);
6920 M.getContext().diagnose(DiagVersion);
6930 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6933 if (
F->hasFnAttribute(Attr)) {
6936 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6938 auto [Part, Rest] = S.
split(
',');
6944 const unsigned Dim = DimC -
'x';
6945 assert(Dim < 3 &&
"Unexpected dim char");
6955 F->addFnAttr(Attr, NewAttr);
6959 return S ==
"x" || S ==
"y" || S ==
"z";
6964 if (
K ==
"kernel") {
6976 const unsigned Idx = (AlignIdxValuePair >> 16);
6977 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6982 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6987 if (
K ==
"minctasm") {
6992 if (
K ==
"maxnreg") {
6997 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
7001 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
7005 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
7009 if (
K ==
"grid_constant") {
7024 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
7031 if (!SeenNodes.
insert(MD).second)
7038 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
7045 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
7047 const MDOperand &V = MD->getOperand(j + 1);
7053 if (NewOperands.
size() > 1)
7066 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
7067 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
7068 if (ModRetainReleaseMarker) {
7074 ID->getString().split(ValueComp,
"#");
7075 if (ValueComp.
size() == 2) {
7076 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
7080 M.eraseNamedMetadata(ModRetainReleaseMarker);
7091 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
7117 bool InvalidCast =
false;
7119 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7132 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7134 Args.push_back(Arg);
7141 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7146 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7159 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7167 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7168 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7169 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7170 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7171 {
"objc_autoreleaseReturnValue",
7172 llvm::Intrinsic::objc_autoreleaseReturnValue},
7173 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7174 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7175 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7176 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7177 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7178 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7179 {
"objc_release", llvm::Intrinsic::objc_release},
7180 {
"objc_retain", llvm::Intrinsic::objc_retain},
7181 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7182 {
"objc_retainAutoreleaseReturnValue",
7183 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7184 {
"objc_retainAutoreleasedReturnValue",
7185 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7186 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7187 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7188 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7189 {
"objc_unsafeClaimAutoreleasedReturnValue",
7190 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7191 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7192 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7193 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7194 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7195 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7196 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7197 {
"objc_arc_annotation_topdown_bbstart",
7198 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7199 {
"objc_arc_annotation_topdown_bbend",
7200 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7201 {
"objc_arc_annotation_bottomup_bbstart",
7202 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7203 {
"objc_arc_annotation_bottomup_bbend",
7204 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7206 for (
auto &
I : RuntimeFuncs)
7207 UpgradeToIntrinsic(
I.first,
I.second);
7231 std::optional<bool> UseAddressDisc;
7234 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7235 for (
const MDNode *Flag : ModFlags->operands()) {
7237 if (Name && (*Name ==
"ptrauth-init-fini" ||
7238 *Name ==
"ptrauth-init-fini-address-discrimination"))
7243 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7244 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7245 constexpr unsigned ExpectedAddressMarker = 1;
7248 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7251 bool HasAddressDisc;
7252 if (!CPA->hasAddressDiscriminator())
7253 HasAddressDisc =
false;
7254 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7255 HasAddressDisc =
true;
7259 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7262 UseAddressDisc = HasAddressDisc;
7263 return CPA->getPointer();
7267 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7270 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7272 if (!GV || !GV->hasInitializer())
7276 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7279 std::vector<Constant *> NewStructors;
7280 NewStructors.reserve(OldStructorsArray->getNumOperands());
7282 for (
Use &U : OldStructorsArray->operands()) {
7291 Func = UpgradeSinglePointer(Func);
7295 NewStructors.push_back(
7304 if (GlobalArraysToUpgrade.
empty())
7306 assert(UseAddressDisc.has_value());
7308 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7309 GV->setInitializer(NewInit);
7312 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7322 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7326 bool HasObjCFlag =
false, HasClassProperties =
false;
7327 bool HasSwiftVersionFlag =
false;
7328 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7335 if (
Op->getNumOperands() != 3)
7349 if (ID->getString() ==
"Objective-C Image Info Version")
7351 if (ID->getString() ==
"Objective-C Class Properties")
7352 HasClassProperties =
true;
7354 if (ID->getString() ==
"PIC Level") {
7355 if (
auto *Behavior =
7357 uint64_t V = Behavior->getLimitedValue();
7363 if (ID->getString() ==
"PIE Level")
7364 if (
auto *Behavior =
7372 if (ID->getString() ==
"branch-target-enforcement" ||
7373 (ID->getString().starts_with(
"sign-return-address") &&
7374 ID->getString() !=
"sign-return-address-harden")) {
7375 if (
auto *Behavior =
7381 Op->getOperand(1),
Op->getOperand(2)};
7391 if (ID->getString() ==
"Objective-C Image Info Section") {
7394 Value->getString().split(ValueComp,
" ");
7395 if (ValueComp.
size() != 1) {
7396 std::string NewValue;
7397 for (
auto &S : ValueComp)
7398 NewValue += S.str();
7409 if (ID->getString() ==
"Objective-C Garbage Collection") {
7412 assert(Md->getValue() &&
"Expected non-empty metadata");
7413 auto Type = Md->getValue()->getType();
7416 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7417 if ((Val & 0xff) != Val) {
7418 HasSwiftVersionFlag =
true;
7419 SwiftABIVersion = (Val & 0xff00) >> 8;
7420 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7421 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7432 if (ID->getString() ==
"amdgpu_code_object_version") {
7435 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7444 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7473 if (HasObjCFlag && !HasClassProperties) {
7479 if (HasSwiftVersionFlag) {
7483 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7485 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7493 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7497 auto MatchesVersion = [](
const MDNode *
Op) {
7498 return Op->getNumOperands() >= 3 &&
7512 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7513 assert(
Op->getNumOperands() >= 2 &&
7514 "Expected at least 2 operands - name and linkage type");
7526 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7537 auto TrimSpaces = [](
StringRef Section) -> std::string {
7539 Section.split(Components,
',');
7544 for (
auto Component : Components)
7545 OS <<
',' << Component.trim();
7550 for (
auto &GV : M.globals()) {
7551 if (!GV.hasSection())
7556 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7561 GV.setSection(TrimSpaces(Section));
7577struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7578 StrictFPUpgradeVisitor() =
default;
7581 if (!
Call.isStrictFP())
7587 Call.removeFnAttr(Attribute::StrictFP);
7588 Call.addFnAttr(Attribute::NoBuiltin);
7593struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7594 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7595 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7597 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7612 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7613 StrictFPUpgradeVisitor SFPV;
7618 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7619 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7620 for (
auto &Arg :
F.args())
7622 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7624 bool AddingAttrs =
false, RemovingAttrs =
false;
7625 AttrBuilder AttrsToAdd(
F.getContext());
7630 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7631 A.isValid() &&
A.isStringAttribute()) {
7632 F.setSection(
A.getValueAsString());
7634 RemovingAttrs =
true;
7638 A.isValid() &&
A.isStringAttribute()) {
7641 AddingAttrs = RemovingAttrs =
true;
7644 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7645 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7647 RemovingAttrs =
true;
7648 if (
A.getValueAsString() ==
"true") {
7649 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7658 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7661 if (
A.getValueAsBool()) {
7662 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7668 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7669 RemovingAttrs =
true;
7676 bool HandleDenormalMode =
false;
7678 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7681 DenormalFPMath = ParsedMode;
7683 AddingAttrs = RemovingAttrs =
true;
7684 HandleDenormalMode =
true;
7688 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7692 DenormalFPMathF32 = ParsedMode;
7694 AddingAttrs = RemovingAttrs =
true;
7695 HandleDenormalMode =
true;
7699 if (HandleDenormalMode)
7700 AttrsToAdd.addDenormalFPEnvAttr(
7704 F.removeFnAttrs(AttrsToRemove);
7707 F.addFnAttrs(AttrsToAdd);
7713 if (!
F.hasFnAttribute(FnAttrName)) {
7714 F.addFnAttr(FnAttrName,
Value);
7716 <<
"\", function: " <<
F.getName() <<
"\n");
7724 if (!
F.hasFnAttribute(FnAttrName)) {
7726 F.addFnAttr(FnAttrName);
7728 <<
", function: " <<
F.getName() <<
"\n");
7731 auto A =
F.getFnAttribute(FnAttrName);
7732 if (
"false" ==
A.getValueAsString()) {
7733 F.removeFnAttr(FnAttrName);
7735 <<
"=\"false\", function: " <<
F.getName() <<
"\n");
7736 }
else if (
"true" ==
A.getValueAsString()) {
7737 F.removeFnAttr(FnAttrName);
7738 F.addFnAttr(FnAttrName);
7740 <<
"=\"true\", function: " <<
F.getName() <<
"\n");
7747 M.setModuleFlag(Behavior,
Key, Val);
7748 LLVM_DEBUG(
dbgs() <<
"Converted module flag: " <<
"{" << Behavior <<
", "
7749 <<
Key <<
", " << Val <<
"}\n");
7753 Triple T(M.getTargetTriple());
7754 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7757 uint64_t BTEValue = 0;
7758 uint64_t BPPLRValue = 0;
7759 uint64_t GCSValue = 0;
7760 uint64_t SRAValue = 0;
7761 uint64_t SRAALLValue = 0;
7762 uint64_t SRABKeyValue = 0;
7764 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7768 if (
Op->getNumOperands() != 3)
7777 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7778 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7779 : IDStr ==
"guarded-control-stack" ? &GCSValue
7780 : IDStr ==
"sign-return-address" ? &SRAValue
7781 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7782 : IDStr ==
"sign-return-address-with-bkey"
7788 *ValPtr = CI->getZExtValue();
7792 LLVM_DEBUG(
dbgs() <<
"Found module flag: " << IDStr <<
"(" << *ValPtr
7797 bool BTE = BTEValue == 1;
7798 bool BPPLR = BPPLRValue == 1;
7799 bool GCS = GCSValue == 1;
7800 bool SRA = SRAValue == 1;
7803 if (SRA && SRAALLValue == 1)
7804 SignTypeValue =
"all";
7807 if (SRA && SRABKeyValue == 1)
7808 SignKeyValue =
"b_key";
7810 for (
Function &
F : M.getFunctionList()) {
7811 if (
F.isDeclaration())
7818 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7819 A.isValid() &&
"none" ==
A.getValueAsString()) {
7820 F.removeFnAttr(
"sign-return-address");
7821 F.removeFnAttr(
"sign-return-address-key");
7837 if (SRAALLValue == 1)
7839 if (SRABKeyValue == 1) {
7868 if (
T->getNumOperands() < 1)
7873 if (S->getString().starts_with(
"llvm.vectorizer."))
7879 StringRef OldPrefix =
"llvm.vectorizer.";
7882 if (OldTag ==
"llvm.vectorizer.unroll")
7894 if (
T->getNumOperands() < 1)
7906 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7919 Ops.reserve(
T->getNumOperands());
7920 Ops.push_back(NewTag);
7921 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7922 Ops.push_back(
T->getOperand(
I));
7939 if (
T->isDistinct()) {
7940 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7952 Ops.reserve(
T->getNumOperands());
7963 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7964 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7965 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7968 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7970 auto I =
DL.find(
"-n64-");
7972 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7977 std::string Res =
DL.str();
7980 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7981 Res.append(Res.empty() ?
"G1" :
"-G1");
7989 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7990 Res.append(
"-ni:7:8:9");
7992 if (
DL.ends_with(
"ni:7"))
7994 if (
DL.ends_with(
"ni:7:8"))
7999 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
8000 Res.append(
"-p7:160:256:256:32");
8001 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
8002 Res.append(
"-p8:128:128:128:48");
8003 constexpr StringRef OldP8(
"-p8:128:128-");
8004 if (
DL.contains(OldP8))
8005 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
8006 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
8007 Res.append(
"-p9:192:256:256:32");
8012 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
8013 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
8014 Res.append((
"-" + AS +
":32:32").str());
8019 if (!
DL.contains(
"m:e"))
8020 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
8025 if (
T.isSystemZ() && !
DL.empty()) {
8027 if (!
DL.contains(
"-S64"))
8028 return "E-S64" +
DL.drop_front(1).str();
8032 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
8035 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
8036 if (!
DL.contains(AddrSpaces)) {
8038 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
8039 if (R.match(Res, &
Groups))
8045 if (
T.isAArch64()) {
8047 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
8048 Res.append(
"-Fn32");
8049 AddPtr32Ptr64AddrSpaces();
8053 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
8057 std::string I64 =
"-i64:64";
8058 std::string I128 =
"-i128:128";
8060 size_t Pos = Res.find(I64);
8061 if (Pos !=
size_t(-1))
8062 Res.insert(Pos + I64.size(), I128);
8066 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
8067 size_t Pos = Res.find(
"-S128");
8070 Res.insert(Pos,
"-f64:32:64");
8075 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
8077 size_t Pos = Res.
find(p3232);
8079 Res.insert(Pos + p3232.
size(),
"-Fi8");
8085 AddPtr32Ptr64AddrSpaces();
8093 if (!
T.isOSIAMCU()) {
8094 std::string I128 =
"-i128:128";
8097 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
8098 if (R.match(Res, &
Groups))
8106 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
8108 auto I =
Ref.find(
"-f80:32-");
8110 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
8118 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
8121 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
8122 B.removeAttribute(
"no-frame-pointer-elim");
8124 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
8126 if (FramePointer !=
"all")
8127 FramePointer =
"non-leaf";
8128 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
8130 if (!FramePointer.
empty())
8131 B.addAttribute(
"frame-pointer", FramePointer);
8133 A =
B.getAttribute(
"null-pointer-is-valid");
8136 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
8137 B.removeAttribute(
"null-pointer-is-valid");
8138 if (NullPointerIsValid)
8139 B.addAttribute(Attribute::NullPointerIsValid);
8142 A =
B.getAttribute(
"uniform-work-group-size");
8146 bool IsTrue = Val ==
"true";
8147 B.removeAttribute(
"uniform-work-group-size");
8149 B.addAttribute(
"uniform-work-group-size");
8160 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
static unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID, SmallVectorImpl< Type * > &OverloadTys)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static bool isValidLegacyNVVMFPArithIntrinsic(const Intrinsic::ID(&IIDs)[2][3], StringRef Type, bool IsRN, bool IsFTZ, bool IsSat, bool IsRelu)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static constexpr Intrinsic::ID NVVMFMAIIDs[2][3]
static constexpr Intrinsic::ID NVVMFMulIIDs[2][3]
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVVMBarrierMulticastIntrinsic(Function *F, StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static void ConvertModuleFlag(Module &M, Module::ModFlagBehavior Behavior, StringRef Key, uint32_t Val)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static Value * upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI, StringRef Name, const Intrinsic::ID(&IIDs)[2][3])
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(Function *F, StringRef Name)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFPArithUpgrade(StringRef Name, const Intrinsic::ID(&IIDs)[2][3])
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F, StringRef Name)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_CASE(ID_SUFFIX, NAME)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static FCmpInst::Predicate getVPFPPredicateFromMD(const Value *Op)
static constexpr Intrinsic::ID NVVMFAddIIDs[2][3]
static GCRegistry::Add< ShadowStackGC > C("shadow-stack", "Very portable GC for uncooperative code generators")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
This file contains constants used for implementing Dwarf debug support.
Module.h This file contains the declarations for the Module class.
const AbstractManglingParser< Derived, Alloc >::OperatorInfo AbstractManglingParser< Derived, Alloc >::Ops[]
static bool isZero(Value *V, const DataLayout &DL, DominatorTree *DT, AssumptionCache *AC)
NVPTX address space definition.
This file contains the definitions of the enumerations and flags associated with NVVM Intrinsics,...
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
static SymbolRef::Type getType(const Symbol *Sym)
LocallyHashedType DenseMapInfo< LocallyHashedType >::Empty
static const X86InstrFMA3Group Groups[]
Class for arbitrary precision integers.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
Class to represent array types.
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Type * getElementType() const
an instruction that atomically reads a memory location, combines it with another value,...
void setVolatile(bool V)
Specify whether this is a volatile RMW or not.
BinOp
This enumeration lists the possible modifications atomicrmw can make.
@ USubCond
Subtract only if no unsigned overflow.
@ Min
*p = old <signed v ? old : v
@ USubSat
*p = usub.sat(old, v) usub.sat matches the behavior of llvm.usub.sat.
@ UIncWrap
Increment one up to a maximum value.
@ Max
*p = old >signed v ? old : v
@ FMin
*p = minnum(old, v) minnum matches the behavior of llvm.minnum.
@ FMax
*p = maxnum(old, v) maxnum matches the behavior of llvm.maxnum.
@ UDecWrap
Decrement one until a minimum value or zero.
bool isFloatingPointOperation() const
This class stores enough information to efficiently remove some attributes from an existing AttrBuild...
AttributeMask & addAttribute(Attribute::AttrKind Val)
Add an attribute to the mask.
Functions, function parameters, and return types can have attributes to indicate how they should be t...
static LLVM_ABI Attribute getWithStackAlignment(LLVMContext &Context, Align Alignment)
static LLVM_ABI Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
void setCallingConv(CallingConv::ID CC)
LLVM_ABI void getOperandBundlesAsDefs(SmallVectorImpl< OperandBundleDef > &Defs) const
Return the list of operand bundles attached to this instruction as a vector of OperandBundleDefs.
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
CallingConv::ID getCallingConv() const
Value * getCalledOperand() const
void setAttributes(AttributeList A)
Set the attributes for this call.
Value * getArgOperand(unsigned i) const
FunctionType * getFunctionType() const
LLVM_ABI Intrinsic::ID getIntrinsicID() const
Returns the intrinsic ID of the intrinsic called or Intrinsic::not_intrinsic if the called function i...
iterator_range< User::op_iterator > args()
Iteration adapter for range-for loops.
void setCalledOperand(Value *V)
unsigned arg_size() const
AttributeList getAttributes() const
Return the attributes for this call.
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
This class represents a function call, abstracting a target machine's calling convention.
void setTailCallKind(TailCallKind TCK)
static LLVM_ABI CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", InsertPosition InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
static LLVM_ABI bool castIsValid(Instruction::CastOps op, Type *SrcTy, Type *DstTy)
This method can be used to determine if a cast from SrcTy to DstTy using Opcode op is valid or not.
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
unsigned getNumParams() const
Return the number of fixed parameters this function type requires.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Module * getParent()
Get the module that this global value is contained inside of...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
op_iterator op_end() const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
op_iterator op_begin() const
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
bool consume_front(char Prefix)
Returns true if this StringRef has the given prefix and removes that prefix.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
StringSwitch & StartsWith(StringLiteral S, T Value)
StringSwitch & Cases(std::initializer_list< StringLiteral > CaseStrings, T Value)
Class to represent struct types.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
unsigned getNumElements() const
Random access to the elements.
Type * getElementType(unsigned N) const
The TimeTraceScope is a helper class to call the begin and end functions of the time trace profiler.
Triple - Helper class for working with autoconf configuration names.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)
bool isVectorTy() const
True if this is an instance of VectorType.
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
LLVM_ABI unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)
LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isPtrOrPtrVectorTy() const
Return true if this is a pointer type or a vector of pointer types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isFPOrFPVectorTy() const
Return true if this is a FP type or a vector of FP.
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
static LLVM_ABI Type * getBFloatTy(LLVMContext &C)
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
A Use represents the edge between a Value definition and its users.
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
LLVM_ABI void print(raw_ostream &O, bool IsForDebug=false) const
Implement operator<< on Value.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set a particular kind of metadata attachment.
LLVM_ABI void setName(const Twine &Name)
Change the name of the value.
LLVM_ABI void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
LLVMContext & getContext() const
All values hold a context through their type.
iterator_range< user_iterator > users()
LLVM_ABI const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
LLVM_ABI void takeName(Value *V)
Transfer the name from V to this value.
Base class of all SIMD vector types.
static VectorType * getInteger(VectorType *VTy)
This static method gets a VectorType with the same number of elements as the input type,...
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
constexpr ScalarTy getFixedValue() const
const ParentTy * getParent() const
self_iterator getIterator()
A raw_ostream that writes to an SmallVector or SmallString.
StringRef str() const
Return a StringRef for the vector contents.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ LOCAL_ADDRESS
Address space for local memory.
@ FLAT_ADDRESS
Address space for flat memory.
@ PRIVATE_ADDRESS
Address space for private memory.
@ PTX_Kernel
Call to a PTX kernel. Passes all arguments in parameter space.
std::optional< ABIType > parseABIType(StringRef S)
Parse the string spelling used by the "float-abi" IR module flag into an ABIType.
LLVM_ABI std::optional< Function * > remangleIntrinsicFunction(Function *F)
LLVM_ABI Function * getOrInsertDeclaration(Module *M, ID id, ArrayRef< Type * > OverloadTys={})
Look up the Function declaration of the intrinsic id in the Module M.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
LLVM_ABI MDNode * UpgradeTBAAStructNode(MDNode &TBAAStructNode)
If the given !tbaa.struct node has old-style scalar field tags, return an equivalent node with each f...
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ATTRIBUTE_VISIBILITY_DEFAULT AnalysisKey InnerAnalysisManagerProxy< AnalysisManagerT, IRUnitT, ExtraArgTs... >::Key
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
const IROptions & getIROptions(const LLVMContext &C)
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
LLVM_ABI bool verifyModule(const Module &M, raw_ostream *OS=nullptr, bool *BrokenDebugInfo=nullptr)
Check a module for errors.
LLVM_ABI void reportFatalUsageError(Error Err)
Report a fatal error that does not indicate a bug in LLVM.
void swap(llvm::BitVector &LHS, llvm::BitVector &RHS)
Implement std::swap in terms of BitVector swap.
This struct is a compact representation of a valid (non-zero power of two) alignment.
Represents the full denormal controls for a function, including the default mode and the f32 specific...
Represent subnormal handling kind for floating point instruction inputs and outputs.
static constexpr DenormalMode getInvalid()
constexpr bool isValid() const
static constexpr DenormalMode getIEEE()
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.