35#include "llvm/IR/IntrinsicsAArch64.h"
36#include "llvm/IR/IntrinsicsAMDGPU.h"
37#include "llvm/IR/IntrinsicsARM.h"
38#include "llvm/IR/IntrinsicsNVPTX.h"
39#include "llvm/IR/IntrinsicsRISCV.h"
40#include "llvm/IR/IntrinsicsWebAssembly.h"
41#include "llvm/IR/IntrinsicsX86.h"
65 cl::desc(
"Disable autoupgrade of debug info"));
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") ||
215 Name ==
"pmulu.dq" ||
216 Name.starts_with(
"psll.dq") ||
217 Name.starts_with(
"psrl.dq") ||
218 Name.starts_with(
"psubs.") ||
219 Name.starts_with(
"psubus.") ||
220 Name.starts_with(
"vbroadcast") ||
221 Name ==
"vbroadcasti128" ||
222 Name ==
"vextracti128" ||
223 Name ==
"vinserti128" ||
224 Name ==
"vperm2i128");
226 if (Name.consume_front(
"avx512.")) {
227 if (Name.consume_front(
"mask."))
229 return (Name.starts_with(
"add.p") ||
230 Name.starts_with(
"and.") ||
231 Name.starts_with(
"andn.") ||
232 Name.starts_with(
"broadcast.s") ||
233 Name.starts_with(
"broadcastf32x4.") ||
234 Name.starts_with(
"broadcastf32x8.") ||
235 Name.starts_with(
"broadcastf64x2.") ||
236 Name.starts_with(
"broadcastf64x4.") ||
237 Name.starts_with(
"broadcasti32x4.") ||
238 Name.starts_with(
"broadcasti32x8.") ||
239 Name.starts_with(
"broadcasti64x2.") ||
240 Name.starts_with(
"broadcasti64x4.") ||
241 Name.starts_with(
"cmp.b") ||
242 Name.starts_with(
"cmp.d") ||
243 Name.starts_with(
"cmp.q") ||
244 Name.starts_with(
"cmp.w") ||
245 Name.starts_with(
"compress.b") ||
246 Name.starts_with(
"compress.d") ||
247 Name.starts_with(
"compress.p") ||
248 Name.starts_with(
"compress.q") ||
249 Name.starts_with(
"compress.store.") ||
250 Name.starts_with(
"compress.w") ||
251 Name.starts_with(
"conflict.") ||
252 Name.starts_with(
"cvtdq2pd.") ||
253 Name.starts_with(
"cvtdq2ps.") ||
254 Name ==
"cvtpd2dq.256" ||
255 Name ==
"cvtpd2ps.256" ||
256 Name ==
"cvtps2pd.128" ||
257 Name ==
"cvtps2pd.256" ||
258 Name.starts_with(
"cvtqq2pd.") ||
259 Name ==
"cvtqq2ps.256" ||
260 Name ==
"cvtqq2ps.512" ||
261 Name ==
"cvttpd2dq.256" ||
262 Name ==
"cvttps2dq.128" ||
263 Name ==
"cvttps2dq.256" ||
264 Name.starts_with(
"cvtudq2pd.") ||
265 Name.starts_with(
"cvtudq2ps.") ||
266 Name.starts_with(
"cvtuqq2pd.") ||
267 Name ==
"cvtuqq2ps.256" ||
268 Name ==
"cvtuqq2ps.512" ||
269 Name.starts_with(
"dbpsadbw.") ||
270 Name.starts_with(
"div.p") ||
271 Name.starts_with(
"expand.b") ||
272 Name.starts_with(
"expand.d") ||
273 Name.starts_with(
"expand.load.") ||
274 Name.starts_with(
"expand.p") ||
275 Name.starts_with(
"expand.q") ||
276 Name.starts_with(
"expand.w") ||
277 Name.starts_with(
"fpclass.p") ||
278 Name.starts_with(
"insert") ||
279 Name.starts_with(
"load.") ||
280 Name.starts_with(
"loadu.") ||
281 Name.starts_with(
"lzcnt.") ||
282 Name.starts_with(
"max.p") ||
283 Name.starts_with(
"min.p") ||
284 Name.starts_with(
"movddup") ||
285 Name.starts_with(
"move.s") ||
286 Name.starts_with(
"movshdup") ||
287 Name.starts_with(
"movsldup") ||
288 Name.starts_with(
"mul.p") ||
289 Name.starts_with(
"or.") ||
290 Name.starts_with(
"pabs.") ||
291 Name.starts_with(
"packssdw.") ||
292 Name.starts_with(
"packsswb.") ||
293 Name.starts_with(
"packusdw.") ||
294 Name.starts_with(
"packuswb.") ||
295 Name.starts_with(
"padd.") ||
296 Name.starts_with(
"padds.") ||
297 Name.starts_with(
"paddus.") ||
298 Name.starts_with(
"palignr.") ||
299 Name.starts_with(
"pand.") ||
300 Name.starts_with(
"pandn.") ||
301 Name.starts_with(
"pavg") ||
302 Name.starts_with(
"pbroadcast") ||
303 Name.starts_with(
"pcmpeq.") ||
304 Name.starts_with(
"pcmpgt.") ||
305 Name.starts_with(
"perm.df.") ||
306 Name.starts_with(
"perm.di.") ||
307 Name.starts_with(
"permvar.") ||
308 Name.starts_with(
"pmaddubs.w.") ||
309 Name.starts_with(
"pmaddw.d.") ||
310 Name.starts_with(
"pmax") ||
311 Name.starts_with(
"pmin") ||
312 Name ==
"pmov.qd.256" ||
313 Name ==
"pmov.qd.512" ||
314 Name ==
"pmov.wb.256" ||
315 Name ==
"pmov.wb.512" ||
316 Name.starts_with(
"pmovsx") ||
317 Name.starts_with(
"pmovzx") ||
318 Name.starts_with(
"pmul.dq.") ||
319 Name.starts_with(
"pmul.hr.sw.") ||
320 Name.starts_with(
"pmulh.w.") ||
321 Name.starts_with(
"pmulhu.w.") ||
322 Name.starts_with(
"pmull.") ||
323 Name.starts_with(
"pmultishift.qb.") ||
324 Name.starts_with(
"pmulu.dq.") ||
325 Name.starts_with(
"por.") ||
326 Name.starts_with(
"prol.") ||
327 Name.starts_with(
"prolv.") ||
328 Name.starts_with(
"pror.") ||
329 Name.starts_with(
"prorv.") ||
330 Name.starts_with(
"pshuf.b.") ||
331 Name.starts_with(
"pshuf.d.") ||
332 Name.starts_with(
"pshufh.w.") ||
333 Name.starts_with(
"pshufl.w.") ||
334 Name.starts_with(
"psll.d") ||
335 Name.starts_with(
"psll.q") ||
336 Name.starts_with(
"psll.w") ||
337 Name.starts_with(
"pslli") ||
338 Name.starts_with(
"psllv") ||
339 Name.starts_with(
"psra.d") ||
340 Name.starts_with(
"psra.q") ||
341 Name.starts_with(
"psra.w") ||
342 Name.starts_with(
"psrai") ||
343 Name.starts_with(
"psrav") ||
344 Name.starts_with(
"psrl.d") ||
345 Name.starts_with(
"psrl.q") ||
346 Name.starts_with(
"psrl.w") ||
347 Name.starts_with(
"psrli") ||
348 Name.starts_with(
"psrlv") ||
349 Name.starts_with(
"psub.") ||
350 Name.starts_with(
"psubs.") ||
351 Name.starts_with(
"psubus.") ||
352 Name.starts_with(
"pternlog.") ||
353 Name.starts_with(
"punpckh") ||
354 Name.starts_with(
"punpckl") ||
355 Name.starts_with(
"pxor.") ||
356 Name.starts_with(
"shuf.f") ||
357 Name.starts_with(
"shuf.i") ||
358 Name.starts_with(
"shuf.p") ||
359 Name.starts_with(
"sqrt.p") ||
360 Name.starts_with(
"store.b.") ||
361 Name.starts_with(
"store.d.") ||
362 Name.starts_with(
"store.p") ||
363 Name.starts_with(
"store.q.") ||
364 Name.starts_with(
"store.w.") ||
365 Name ==
"store.ss" ||
366 Name.starts_with(
"storeu.") ||
367 Name.starts_with(
"sub.p") ||
368 Name.starts_with(
"ucmp.") ||
369 Name.starts_with(
"unpckh.") ||
370 Name.starts_with(
"unpckl.") ||
371 Name.starts_with(
"valign.") ||
372 Name ==
"vcvtph2ps.128" ||
373 Name ==
"vcvtph2ps.256" ||
374 Name.starts_with(
"vextract") ||
375 Name.starts_with(
"vfmadd.") ||
376 Name.starts_with(
"vfmaddsub.") ||
377 Name.starts_with(
"vfnmadd.") ||
378 Name.starts_with(
"vfnmsub.") ||
379 Name.starts_with(
"vpdpbusd.") ||
380 Name.starts_with(
"vpdpbusds.") ||
381 Name.starts_with(
"vpdpwssd.") ||
382 Name.starts_with(
"vpdpwssds.") ||
383 Name.starts_with(
"vpermi2var.") ||
384 Name.starts_with(
"vpermil.p") ||
385 Name.starts_with(
"vpermilvar.") ||
386 Name.starts_with(
"vpermt2var.") ||
387 Name.starts_with(
"vpmadd52") ||
388 Name.starts_with(
"vpshld.") ||
389 Name.starts_with(
"vpshldv.") ||
390 Name.starts_with(
"vpshrd.") ||
391 Name.starts_with(
"vpshrdv.") ||
392 Name.starts_with(
"vpshufbitqmb.") ||
393 Name.starts_with(
"xor."));
395 if (Name.consume_front(
"mask3."))
397 return (Name.starts_with(
"vfmadd.") ||
398 Name.starts_with(
"vfmaddsub.") ||
399 Name.starts_with(
"vfmsub.") ||
400 Name.starts_with(
"vfmsubadd.") ||
401 Name.starts_with(
"vfnmsub."));
403 if (Name.consume_front(
"maskz."))
405 return (Name.starts_with(
"pternlog.") ||
406 Name.starts_with(
"vfmadd.") ||
407 Name.starts_with(
"vfmaddsub.") ||
408 Name.starts_with(
"vpdpbusd.") ||
409 Name.starts_with(
"vpdpbusds.") ||
410 Name.starts_with(
"vpdpwssd.") ||
411 Name.starts_with(
"vpdpwssds.") ||
412 Name.starts_with(
"vpermt2var.") ||
413 Name.starts_with(
"vpmadd52") ||
414 Name.starts_with(
"vpshldv.") ||
415 Name.starts_with(
"vpshrdv."));
418 return (Name ==
"movntdqa" ||
419 Name ==
"pmul.dq.512" ||
420 Name ==
"pmulu.dq.512" ||
421 Name.starts_with(
"broadcastm") ||
422 Name.starts_with(
"cmp.p") ||
423 Name.starts_with(
"cvtb2mask.") ||
424 Name.starts_with(
"cvtd2mask.") ||
425 Name.starts_with(
"cvtmask2") ||
426 Name.starts_with(
"cvtq2mask.") ||
427 Name ==
"cvtusi2sd" ||
428 Name.starts_with(
"cvtw2mask.") ||
433 Name ==
"kortestc.w" ||
434 Name ==
"kortestz.w" ||
435 Name.starts_with(
"kunpck") ||
438 Name.starts_with(
"padds.") ||
439 Name.starts_with(
"pbroadcast") ||
440 Name.starts_with(
"prol") ||
441 Name.starts_with(
"pror") ||
442 Name.starts_with(
"psll.dq") ||
443 Name.starts_with(
"psrl.dq") ||
444 Name.starts_with(
"psubs.") ||
445 Name.starts_with(
"ptestm") ||
446 Name.starts_with(
"ptestnm") ||
447 Name.starts_with(
"storent.") ||
448 Name.starts_with(
"vbroadcast.s") ||
449 Name.starts_with(
"vpshld.") ||
450 Name.starts_with(
"vpshrd."));
453 if (Name.consume_front(
"fma."))
454 return (Name.starts_with(
"vfmadd.") ||
455 Name.starts_with(
"vfmsub.") ||
456 Name.starts_with(
"vfmsubadd.") ||
457 Name.starts_with(
"vfnmadd.") ||
458 Name.starts_with(
"vfnmsub."));
460 if (Name.consume_front(
"fma4."))
461 return Name.starts_with(
"vfmadd.s");
463 if (Name.consume_front(
"sse."))
464 return (Name ==
"add.ss" ||
465 Name ==
"cvtsi2ss" ||
466 Name ==
"cvtsi642ss" ||
469 Name.starts_with(
"sqrt.p") ||
471 Name.starts_with(
"storeu.") ||
474 if (Name.consume_front(
"sse2."))
475 return (Name ==
"add.sd" ||
476 Name ==
"cvtdq2pd" ||
477 Name ==
"cvtdq2ps" ||
478 Name ==
"cvtps2pd" ||
479 Name ==
"cvtsi2sd" ||
480 Name ==
"cvtsi642sd" ||
481 Name ==
"cvtss2sd" ||
484 Name.starts_with(
"padds.") ||
485 Name.starts_with(
"paddus.") ||
486 Name.starts_with(
"pcmpeq.") ||
487 Name.starts_with(
"pcmpgt.") ||
492 Name ==
"pmulu.dq" ||
493 Name.starts_with(
"pshuf") ||
494 Name.starts_with(
"psll.dq") ||
495 Name.starts_with(
"psrl.dq") ||
496 Name.starts_with(
"psubs.") ||
497 Name.starts_with(
"psubus.") ||
498 Name.starts_with(
"sqrt.p") ||
500 Name ==
"storel.dq" ||
501 Name.starts_with(
"storeu.") ||
504 if (Name.consume_front(
"sse41."))
505 return (Name.starts_with(
"blendp") ||
506 Name ==
"movntdqa" ||
516 Name.starts_with(
"pmovsx") ||
517 Name.starts_with(
"pmovzx") ||
520 if (Name.consume_front(
"sse42."))
521 return Name ==
"crc32.64.8";
523 if (Name.consume_front(
"sse4a."))
524 return Name.starts_with(
"movnt.");
526 if (Name.consume_front(
"ssse3."))
527 return (Name ==
"pabs.b.128" ||
528 Name ==
"pabs.d.128" ||
529 Name ==
"pabs.w.128");
531 if (Name.consume_front(
"xop."))
532 return (Name ==
"vpcmov" ||
533 Name ==
"vpcmov.256" ||
534 Name.starts_with(
"vpcom") ||
535 Name.starts_with(
"vprot"));
537 if (Name.consume_front(
"bmi."))
538 return (Name.starts_with(
"pdep.") ||
539 Name.starts_with(
"pext."));
541 return (Name ==
"addcarry.u32" ||
542 Name ==
"addcarry.u64" ||
543 Name ==
"addcarryx.u32" ||
544 Name ==
"addcarryx.u64" ||
545 Name ==
"subborrow.u32" ||
546 Name ==
"subborrow.u64" ||
547 Name.starts_with(
"vcvtph2ps."));
553 if (!Name.consume_front(
"x86."))
561 if (Name ==
"rdtscp") {
563 if (
F->getFunctionType()->getNumParams() == 0)
568 Intrinsic::x86_rdtscp);
575 if (Name.consume_front(
"sse41.ptest")) {
577 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
578 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
579 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
592 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
593 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
594 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
595 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
596 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
597 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
602 if (Name.consume_front(
"avx512.")) {
603 if (Name.consume_front(
"mask.cmp.")) {
606 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
607 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
608 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
609 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
610 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
611 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
615 }
else if (Name.starts_with(
"vpdpbusd.") ||
616 Name.starts_with(
"vpdpbusds.")) {
619 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
620 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
621 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
622 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
623 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
624 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
628 }
else if (Name.starts_with(
"vpdpwssd.") ||
629 Name.starts_with(
"vpdpwssds.")) {
632 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
633 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
634 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
635 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
636 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
637 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
645 if (Name.consume_front(
"avx2.")) {
646 if (Name.consume_front(
"vpdpb")) {
649 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
650 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
651 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
652 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
653 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
654 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
655 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
656 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
657 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
658 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
659 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
660 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
664 }
else if (Name.consume_front(
"vpdpw")) {
667 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
668 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
669 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
670 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
671 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
672 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
673 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
674 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
675 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
676 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
677 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
678 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
686 if (Name.consume_front(
"avx10.")) {
687 if (Name.consume_front(
"vpdpb")) {
690 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
691 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
692 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
693 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
694 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
695 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
699 }
else if (Name.consume_front(
"vpdpw")) {
701 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
702 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
703 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
704 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
705 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
706 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
714 if (Name.consume_front(
"avx512bf16.")) {
717 .
Case(
"cvtne2ps2bf16.128",
718 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
719 .
Case(
"cvtne2ps2bf16.256",
720 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
721 .
Case(
"cvtne2ps2bf16.512",
722 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
723 .
Case(
"mask.cvtneps2bf16.128",
724 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
725 .
Case(
"cvtneps2bf16.256",
726 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
727 .
Case(
"cvtneps2bf16.512",
728 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
735 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
736 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
737 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
744 if (Name.consume_front(
"xop.")) {
746 if (Name.starts_with(
"vpermil2")) {
749 auto Idx =
F->getFunctionType()->getParamType(2);
750 if (Idx->isFPOrFPVectorTy()) {
751 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
752 unsigned EltSize = Idx->getScalarSizeInBits();
753 if (EltSize == 64 && IdxSize == 128)
754 ID = Intrinsic::x86_xop_vpermil2pd;
755 else if (EltSize == 32 && IdxSize == 128)
756 ID = Intrinsic::x86_xop_vpermil2ps;
757 else if (EltSize == 64 && IdxSize == 256)
758 ID = Intrinsic::x86_xop_vpermil2pd_256;
760 ID = Intrinsic::x86_xop_vpermil2ps_256;
762 }
else if (
F->arg_size() == 2)
765 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
766 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
777 if (Name ==
"seh.recoverfp") {
779 Intrinsic::eh_recoverfp);
791 if (Name.starts_with(
"rbit")) {
794 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
798 if (Name ==
"thread.pointer") {
801 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
805 bool Neon = Name.consume_front(
"neon.");
810 if (Name.consume_front(
"bfdot.")) {
814 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
819 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
820 assert((OperandWidth == 64 || OperandWidth == 128) &&
821 "Unexpected operand width");
823 std::array<Type *, 2> Tys{
834 if (Name.consume_front(
"bfm")) {
836 if (Name.consume_back(
".v4f32.v16i8")) {
882 F->arg_begin()->getType());
886 if (Name.consume_front(
"vst")) {
888 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
892 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
893 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
896 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
897 Intrinsic::arm_neon_vst4lane};
899 auto fArgs =
F->getFunctionType()->params();
900 Type *Tys[] = {fArgs[0], fArgs[1]};
903 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
906 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
915 if (Name.consume_front(
"mve.")) {
917 if (Name ==
"vctp64") {
927 if (Name.starts_with(
"vrintn.v")) {
929 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
934 if (Name.consume_back(
".v4i1")) {
936 if (Name.consume_back(
".predicated.v2i64.v4i32"))
938 return Name ==
"mull.int" || Name ==
"vqdmull";
940 if (Name.consume_back(
".v2i64")) {
942 bool IsGather = Name.consume_front(
"vldr.gather.");
943 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
944 if (Name.consume_front(
"base.")) {
946 Name.consume_front(
"wb.");
949 return Name ==
"predicated.v2i64";
952 if (Name.consume_front(
"offset.predicated."))
953 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
954 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
967 if (Name.consume_front(
"cde.vcx")) {
969 if (Name.consume_back(
".predicated.v2i64.v4i1"))
971 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
972 Name ==
"3q" || Name ==
"3qa";
986 F->arg_begin()->getType());
990 if (Name.starts_with(
"addp")) {
992 if (
F->arg_size() != 2)
995 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
997 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1003 if (Name.starts_with(
"bfcvt")) {
1009 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1016 if (Name.consume_front(
"sve.")) {
1018 if (Name.consume_front(
"bf")) {
1019 if (Name ==
"mmla") {
1020 Type *Tys[] = {
F->getReturnType(),
1021 std::next(
F->arg_begin())->getType()};
1023 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1026 if (Name.consume_back(
".lane")) {
1030 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1031 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1032 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1044 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1049 if (Name.consume_front(
"addqv")) {
1051 if (!
F->getReturnType()->isFPOrFPVectorTy())
1054 auto Args =
F->getFunctionType()->params();
1055 Type *Tys[] = {
F->getReturnType(), Args[1]};
1057 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1061 if (Name.consume_front(
"ld")) {
1063 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1064 if (LdRegex.
match(Name)) {
1070 "Expected 2 arguments for ld* intrinsic.");
1071 Type *PtrTy =
F->getArg(1)->getType();
1074 Intrinsic::aarch64_sve_ld2_sret,
1075 Intrinsic::aarch64_sve_ld3_sret,
1076 Intrinsic::aarch64_sve_ld4_sret,
1079 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1085 if (Name.consume_front(
"tuple.")) {
1087 if (Name.starts_with(
"get")) {
1089 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1091 F->getParent(), Intrinsic::vector_extract, Tys);
1095 if (Name.starts_with(
"set")) {
1097 auto Args =
F->getFunctionType()->params();
1098 Type *Tys[] = {Args[0], Args[2], Args[1]};
1100 F->getParent(), Intrinsic::vector_insert, Tys);
1104 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1105 if (CreateTupleRegex.
match(Name)) {
1107 auto Args =
F->getFunctionType()->params();
1108 Type *Tys[] = {
F->getReturnType(), Args[1]};
1110 F->getParent(), Intrinsic::vector_insert, Tys);
1116 if (Name.starts_with(
"rev.nxv")) {
1119 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1125 if (Name.consume_front(
"sme.")) {
1127 if (Name.consume_front(
"ftmopa.")) {
1132 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1133 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1150 if (Name.consume_front(
"cp.async.bulk.tensor.g2s.")) {
1154 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d)
1156 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d)
1158 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d)
1159 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d)
1160 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d)
1161 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d)
1162 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d)
1163 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d)
1172 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1186 size_t FlagStartIndex =
F->getFunctionType()->getNumParams() - 3;
1187 Type *ArgType =
F->getFunctionType()->getParamType(FlagStartIndex);
1197 if (Name.consume_front(
"mapa.shared.cluster"))
1198 if (
F->getReturnType()->getPointerAddressSpace() ==
1200 return Intrinsic::nvvm_mapa_shared_cluster;
1202 if (Name.consume_front(
"cp.async.bulk.")) {
1205 .
Case(
"global.to.shared.cluster",
1206 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1207 .
Case(
"shared.cta.to.cluster",
1208 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1212 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1221 if (Name.consume_front(
"fma.rn."))
1223 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1224 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1225 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1226 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1229 if (Name.consume_front(
"fmax."))
1231 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1232 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1233 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1234 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1235 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1236 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1237 .
Case(
"ftz.nan.xorsign.abs.bf16",
1238 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1239 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1240 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1241 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1242 .
Case(
"ftz.xorsign.abs.bf16x2",
1243 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1244 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1245 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1246 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1247 .
Case(
"nan.xorsign.abs.bf16x2",
1248 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1249 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1250 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1253 if (Name.consume_front(
"fmin."))
1255 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1256 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1257 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1258 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1259 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1260 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1261 .
Case(
"ftz.nan.xorsign.abs.bf16",
1262 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1263 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1264 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1265 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1266 .
Case(
"ftz.xorsign.abs.bf16x2",
1267 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1268 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1269 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1270 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1271 .
Case(
"nan.xorsign.abs.bf16x2",
1272 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1273 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1274 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1277 if (Name.consume_front(
"neg."))
1279 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1280 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1287 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1288 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1289 Name.consume_front(
"param");
1295 if (Name.starts_with(
"to.fp16")) {
1299 FuncTy->getReturnType());
1302 if (Name.starts_with(
"from.fp16")) {
1306 FuncTy->getReturnType());
1318 if (Defaults.empty())
1330 if (
F->arg_size() >= FullDecl->
arg_size())
1335 if (
F->arg_size() < FirstDefault)
1343 bool CanUpgradeDebugIntrinsicsToRecords) {
1344 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1349 if (!Name.consume_front(
"llvm.") || Name.empty())
1355 bool IsArm = Name.consume_front(
"arm.");
1356 if (IsArm || Name.consume_front(
"aarch64.")) {
1362 if (Name.consume_front(
"amdgcn.")) {
1363 if (Name ==
"alignbit") {
1366 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1370 if (Name.consume_front(
"atomic.")) {
1371 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1372 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1381 switch (
F->getIntrinsicID()) {
1385 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1386 if (
F->arg_size() == 7) {
1391 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1392 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1393 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1394 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1395 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1396 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1397 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1398 if (
F->arg_size() == 8) {
1405 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1406 Name.consume_front(
"flat.atomic.")) {
1407 if (Name.starts_with(
"fadd") ||
1409 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1410 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1418 if (Name.starts_with(
"ldexp.")) {
1421 F->getParent(), Intrinsic::ldexp,
1422 {F->getReturnType(), F->getArg(1)->getType()});
1431 if (
F->arg_size() == 1) {
1432 if (Name.consume_front(
"convert.")) {
1446 F->arg_begin()->getType());
1452 if (Name ==
"coro.end" &&
1453 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1454 CoroEndID = Intrinsic::coro_end;
1455 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1456 CoroEndID = Intrinsic::coro_end_async;
1467 if (Name.consume_front(
"dbg.")) {
1469 if (CanUpgradeDebugIntrinsicsToRecords) {
1470 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1471 Name ==
"declare" || Name ==
"label") {
1480 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1483 Intrinsic::dbg_value);
1490 if (Name.consume_front(
"experimental.vector.")) {
1496 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1497 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1498 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1499 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1500 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1502 Intrinsic::vector_partial_reduce_add)
1505 const auto *FT =
F->getFunctionType();
1507 if (ID == Intrinsic::vector_extract ||
1508 ID == Intrinsic::vector_interleave2)
1511 if (ID != Intrinsic::vector_interleave2)
1513 if (ID == Intrinsic::vector_insert ||
1514 ID == Intrinsic::vector_partial_reduce_add)
1522 if (Name.consume_front(
"reduce.")) {
1524 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1525 if (R.match(Name, &
Groups))
1527 .
Case(
"add", Intrinsic::vector_reduce_add)
1528 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1529 .
Case(
"and", Intrinsic::vector_reduce_and)
1530 .
Case(
"or", Intrinsic::vector_reduce_or)
1531 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1532 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1533 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1534 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1535 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1536 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1537 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1542 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1547 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1548 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1553 auto Args =
F->getFunctionType()->params();
1555 {Args[V2 ? 1 : 0]});
1561 if (Name.consume_front(
"splice"))
1565 if (Name.consume_front(
"experimental.stepvector.")) {
1569 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1574 if (Name.starts_with(
"flt.rounds")) {
1577 Intrinsic::get_rounding);
1582 if (Name.starts_with(
"invariant.group.barrier")) {
1584 auto Args =
F->getFunctionType()->params();
1585 Type* ObjectPtr[1] = {Args[0]};
1588 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1593 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1594 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1595 if (IsLifetimeStart || IsLifetimeEnd) {
1596 if (
F->arg_size() == 2) {
1597 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1598 : Intrinsic::lifetime_end;
1603 F->getArg(1)->getType());
1605 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1625 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1626 .StartsWith(
"memmove.", Intrinsic::memmove)
1628 if (
F->arg_size() == 5) {
1632 F->getFunctionType()->params().slice(0, 3);
1638 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1641 const auto *FT =
F->getFunctionType();
1642 Type *ParamTypes[2] = {
1643 FT->getParamType(0),
1647 Intrinsic::memset, ParamTypes);
1653 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1654 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1655 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1656 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1658 if (MaskedID &&
F->arg_size() == 4) {
1660 if (MaskedID == Intrinsic::masked_load ||
1661 MaskedID == Intrinsic::masked_gather) {
1663 F->getParent(), MaskedID,
1664 {F->getReturnType(), F->getArg(0)->getType()});
1668 F->getParent(), MaskedID,
1669 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1675 if (Name.consume_front(
"nvvm.")) {
1677 if (
F->arg_size() == 1) {
1680 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1681 .Case(
"clz.i", Intrinsic::ctlz)
1682 .
Case(
"popc.i", Intrinsic::ctpop)
1686 {F->getReturnType()});
1689 }
else if (
F->arg_size() == 2) {
1692 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1693 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1694 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1695 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1699 {F->getReturnType()});
1705 if (!
F->getReturnType()->getScalarType()->isBFloatTy()) {
1733 bool Expand =
false;
1734 if (Name.consume_front(
"abs."))
1737 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
1738 else if (Name.consume_front(
"fabs."))
1740 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
1741 else if (Name.consume_front(
"ex2.approx."))
1744 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
1745 else if (Name.consume_front(
"atomic.load."))
1754 else if (Name.consume_front(
"atomic."))
1769 else if (Name.consume_front(
"bitcast."))
1772 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
1773 else if (Name.consume_front(
"rotate."))
1775 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
1776 else if (Name.consume_front(
"ptr.gen.to."))
1779 else if (Name.consume_front(
"ptr."))
1782 else if (Name.consume_front(
"ldg.global."))
1784 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
1785 Name.starts_with(
"p."));
1788 .
Case(
"barrier0",
true)
1789 .
Case(
"barrier.n",
true)
1790 .
Case(
"barrier.sync.cnt",
true)
1791 .
Case(
"barrier.sync",
true)
1792 .
Case(
"barrier",
true)
1793 .
Case(
"bar.sync",
true)
1794 .
Case(
"barrier0.popc",
true)
1795 .
Case(
"barrier0.and",
true)
1796 .
Case(
"barrier0.or",
true)
1797 .
Case(
"clz.ll",
true)
1798 .
Case(
"popc.ll",
true)
1800 .
Case(
"swap.lo.hi.b64",
true)
1801 .
Case(
"tanh.approx.f32",
true)
1813 if (Name.starts_with(
"objectsize.")) {
1814 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
1815 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
1818 Intrinsic::objectsize, Tys);
1825 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
1828 F->getParent(), Intrinsic::ptr_annotation,
1829 {F->arg_begin()->getType(), F->getArg(1)->getType()});
1835 if (Name.consume_front(
"riscv.")) {
1838 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
1839 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
1840 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
1841 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
1844 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
1857 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
1858 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
1867 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
1868 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
1869 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
1870 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
1875 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
1884 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
1886 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
1895 if (Name ==
"stackprotectorcheck") {
1902 if (Name ==
"thread.pointer") {
1904 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
1910 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
1913 F->getParent(), Intrinsic::var_annotation,
1914 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
1917 if (Name.consume_front(
"vector.splice")) {
1918 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
1926 if (Name.consume_front(
"wasm.")) {
1929 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
1930 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
1931 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
1936 F->getReturnType());
1940 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
1942 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
1944 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
1963 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
1973 std::string
Name =
F->getName().str();
1976 Name,
F->getParent());
1987 if (Result != std::nullopt) {
2003 bool CanUpgradeDebugIntrinsicsToRecords) {
2023 GV->
getName() ==
"llvm.global_dtors")) ||
2038 unsigned N =
Init->getNumOperands();
2039 std::vector<Constant *> NewCtors(
N);
2040 for (
unsigned i = 0; i !=
N; ++i) {
2043 Ctor->getAggregateElement(1),
2057 unsigned NumElts = ResultTy->getNumElements() * 8;
2061 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2071 for (
unsigned l = 0; l != NumElts; l += 16)
2072 for (
unsigned i = 0; i != 16; ++i) {
2073 unsigned Idx = NumElts + i - Shift;
2075 Idx -= NumElts - 16;
2076 Idxs[l + i] = Idx + l;
2079 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2083 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2091 unsigned NumElts = ResultTy->getNumElements() * 8;
2095 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2105 for (
unsigned l = 0; l != NumElts; l += 16)
2106 for (
unsigned i = 0; i != 16; ++i) {
2107 unsigned Idx = i + Shift;
2109 Idx += NumElts - 16;
2110 Idxs[l + i] = Idx + l;
2113 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2117 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2125 Mask = Builder.CreateBitCast(Mask, MaskTy);
2131 for (
unsigned i = 0; i != NumElts; ++i)
2133 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2144 if (
C->isAllOnesValue())
2149 return Builder.CreateSelect(Mask, Op0, Op1);
2156 if (
C->isAllOnesValue())
2160 Mask->getType()->getIntegerBitWidth());
2161 Mask = Builder.CreateBitCast(Mask, MaskTy);
2162 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2163 return Builder.CreateSelect(Mask, Op0, Op1);
2176 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2177 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2182 ShiftVal &= (NumElts - 1);
2191 if (ShiftVal > 16) {
2199 for (
unsigned l = 0; l < NumElts; l += 16) {
2200 for (
unsigned i = 0; i != 16; ++i) {
2201 unsigned Idx = ShiftVal + i;
2202 if (!IsVALIGN && Idx >= 16)
2203 Idx += NumElts - 16;
2204 Indices[l + i] = Idx + l;
2209 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2215 bool ZeroMask,
bool IndexForm) {
2218 unsigned EltWidth = Ty->getScalarSizeInBits();
2219 bool IsFloat = Ty->isFPOrFPVectorTy();
2221 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2222 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2223 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2224 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2225 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2226 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2227 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2228 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2229 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2230 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2231 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2232 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2233 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2234 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2235 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2236 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2237 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2238 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2239 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2240 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2241 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2242 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2243 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2244 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2245 else if (VecWidth == 128 && EltWidth == 16)
2246 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2247 else if (VecWidth == 256 && EltWidth == 16)
2248 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2249 else if (VecWidth == 512 && EltWidth == 16)
2250 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2251 else if (VecWidth == 128 && EltWidth == 8)
2252 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2253 else if (VecWidth == 256 && EltWidth == 8)
2254 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2255 else if (VecWidth == 512 && EltWidth == 8)
2256 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2267 Value *V = Builder.CreateIntrinsic(IID, Args);
2279 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2290 bool IsRotateRight) {
2300 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2301 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2304 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2305 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2350 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2355 bool IsShiftRight,
bool ZeroMask) {
2369 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2370 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2373 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2374 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2389 const Align Alignment =
2391 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2396 if (
C->isAllOnesValue())
2397 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2402 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2408 const Align Alignment =
2417 if (
C->isAllOnesValue())
2418 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2423 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2429 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2430 {Op0, Builder.getInt1(
false)});
2445 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2446 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2447 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2448 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2449 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2452 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2453 LHS = Builder.CreateAnd(
LHS, Mask);
2454 RHS = Builder.CreateAnd(
RHS, Mask);
2471 if (!
C || !
C->isAllOnesValue())
2472 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2477 for (
unsigned i = 0; i != NumElts; ++i)
2479 for (
unsigned i = NumElts; i != 8; ++i)
2480 Indices[i] = NumElts + i % NumElts;
2481 Vec = Builder.CreateShuffleVector(Vec,
2485 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2489 unsigned CC,
bool Signed) {
2497 }
else if (CC == 7) {
2533 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2534 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2536 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2537 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2546 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2552 Name = Name.substr(12);
2557 if (Name.starts_with(
"max.p")) {
2558 if (VecWidth == 128 && EltWidth == 32)
2559 IID = Intrinsic::x86_sse_max_ps;
2560 else if (VecWidth == 128 && EltWidth == 64)
2561 IID = Intrinsic::x86_sse2_max_pd;
2562 else if (VecWidth == 256 && EltWidth == 32)
2563 IID = Intrinsic::x86_avx_max_ps_256;
2564 else if (VecWidth == 256 && EltWidth == 64)
2565 IID = Intrinsic::x86_avx_max_pd_256;
2568 }
else if (Name.starts_with(
"min.p")) {
2569 if (VecWidth == 128 && EltWidth == 32)
2570 IID = Intrinsic::x86_sse_min_ps;
2571 else if (VecWidth == 128 && EltWidth == 64)
2572 IID = Intrinsic::x86_sse2_min_pd;
2573 else if (VecWidth == 256 && EltWidth == 32)
2574 IID = Intrinsic::x86_avx_min_ps_256;
2575 else if (VecWidth == 256 && EltWidth == 64)
2576 IID = Intrinsic::x86_avx_min_pd_256;
2579 }
else if (Name.starts_with(
"pshuf.b.")) {
2580 if (VecWidth == 128)
2581 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2582 else if (VecWidth == 256)
2583 IID = Intrinsic::x86_avx2_pshuf_b;
2584 else if (VecWidth == 512)
2585 IID = Intrinsic::x86_avx512_pshuf_b_512;
2588 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2589 if (VecWidth == 128)
2590 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2591 else if (VecWidth == 256)
2592 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2593 else if (VecWidth == 512)
2594 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2597 }
else if (Name.starts_with(
"pmulh.w.")) {
2598 if (VecWidth == 128)
2599 IID = Intrinsic::x86_sse2_pmulh_w;
2600 else if (VecWidth == 256)
2601 IID = Intrinsic::x86_avx2_pmulh_w;
2602 else if (VecWidth == 512)
2603 IID = Intrinsic::x86_avx512_pmulh_w_512;
2606 }
else if (Name.starts_with(
"pmulhu.w.")) {
2607 if (VecWidth == 128)
2608 IID = Intrinsic::x86_sse2_pmulhu_w;
2609 else if (VecWidth == 256)
2610 IID = Intrinsic::x86_avx2_pmulhu_w;
2611 else if (VecWidth == 512)
2612 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2615 }
else if (Name.starts_with(
"pmaddw.d.")) {
2616 if (VecWidth == 128)
2617 IID = Intrinsic::x86_sse2_pmadd_wd;
2618 else if (VecWidth == 256)
2619 IID = Intrinsic::x86_avx2_pmadd_wd;
2620 else if (VecWidth == 512)
2621 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2624 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2625 if (VecWidth == 128)
2626 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2627 else if (VecWidth == 256)
2628 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2629 else if (VecWidth == 512)
2630 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2633 }
else if (Name.starts_with(
"packsswb.")) {
2634 if (VecWidth == 128)
2635 IID = Intrinsic::x86_sse2_packsswb_128;
2636 else if (VecWidth == 256)
2637 IID = Intrinsic::x86_avx2_packsswb;
2638 else if (VecWidth == 512)
2639 IID = Intrinsic::x86_avx512_packsswb_512;
2642 }
else if (Name.starts_with(
"packssdw.")) {
2643 if (VecWidth == 128)
2644 IID = Intrinsic::x86_sse2_packssdw_128;
2645 else if (VecWidth == 256)
2646 IID = Intrinsic::x86_avx2_packssdw;
2647 else if (VecWidth == 512)
2648 IID = Intrinsic::x86_avx512_packssdw_512;
2651 }
else if (Name.starts_with(
"packuswb.")) {
2652 if (VecWidth == 128)
2653 IID = Intrinsic::x86_sse2_packuswb_128;
2654 else if (VecWidth == 256)
2655 IID = Intrinsic::x86_avx2_packuswb;
2656 else if (VecWidth == 512)
2657 IID = Intrinsic::x86_avx512_packuswb_512;
2660 }
else if (Name.starts_with(
"packusdw.")) {
2661 if (VecWidth == 128)
2662 IID = Intrinsic::x86_sse41_packusdw;
2663 else if (VecWidth == 256)
2664 IID = Intrinsic::x86_avx2_packusdw;
2665 else if (VecWidth == 512)
2666 IID = Intrinsic::x86_avx512_packusdw_512;
2669 }
else if (Name.starts_with(
"vpermilvar.")) {
2670 if (VecWidth == 128 && EltWidth == 32)
2671 IID = Intrinsic::x86_avx_vpermilvar_ps;
2672 else if (VecWidth == 128 && EltWidth == 64)
2673 IID = Intrinsic::x86_avx_vpermilvar_pd;
2674 else if (VecWidth == 256 && EltWidth == 32)
2675 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2676 else if (VecWidth == 256 && EltWidth == 64)
2677 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2678 else if (VecWidth == 512 && EltWidth == 32)
2679 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2680 else if (VecWidth == 512 && EltWidth == 64)
2681 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2684 }
else if (Name ==
"cvtpd2dq.256") {
2685 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2686 }
else if (Name ==
"cvtpd2ps.256") {
2687 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2688 }
else if (Name ==
"cvttpd2dq.256") {
2689 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2690 }
else if (Name ==
"cvttps2dq.128") {
2691 IID = Intrinsic::x86_sse2_cvttps2dq;
2692 }
else if (Name ==
"cvttps2dq.256") {
2693 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2694 }
else if (Name.starts_with(
"permvar.")) {
2696 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2697 IID = Intrinsic::x86_avx2_permps;
2698 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2699 IID = Intrinsic::x86_avx2_permd;
2700 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2701 IID = Intrinsic::x86_avx512_permvar_df_256;
2702 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2703 IID = Intrinsic::x86_avx512_permvar_di_256;
2704 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2705 IID = Intrinsic::x86_avx512_permvar_sf_512;
2706 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2707 IID = Intrinsic::x86_avx512_permvar_si_512;
2708 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2709 IID = Intrinsic::x86_avx512_permvar_df_512;
2710 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2711 IID = Intrinsic::x86_avx512_permvar_di_512;
2712 else if (VecWidth == 128 && EltWidth == 16)
2713 IID = Intrinsic::x86_avx512_permvar_hi_128;
2714 else if (VecWidth == 256 && EltWidth == 16)
2715 IID = Intrinsic::x86_avx512_permvar_hi_256;
2716 else if (VecWidth == 512 && EltWidth == 16)
2717 IID = Intrinsic::x86_avx512_permvar_hi_512;
2718 else if (VecWidth == 128 && EltWidth == 8)
2719 IID = Intrinsic::x86_avx512_permvar_qi_128;
2720 else if (VecWidth == 256 && EltWidth == 8)
2721 IID = Intrinsic::x86_avx512_permvar_qi_256;
2722 else if (VecWidth == 512 && EltWidth == 8)
2723 IID = Intrinsic::x86_avx512_permvar_qi_512;
2726 }
else if (Name.starts_with(
"dbpsadbw.")) {
2727 if (VecWidth == 128)
2728 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2729 else if (VecWidth == 256)
2730 IID = Intrinsic::x86_avx512_dbpsadbw_256;
2731 else if (VecWidth == 512)
2732 IID = Intrinsic::x86_avx512_dbpsadbw_512;
2735 }
else if (Name.starts_with(
"pmultishift.qb.")) {
2736 if (VecWidth == 128)
2737 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
2738 else if (VecWidth == 256)
2739 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
2740 else if (VecWidth == 512)
2741 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
2744 }
else if (Name.starts_with(
"conflict.")) {
2745 if (Name[9] ==
'd' && VecWidth == 128)
2746 IID = Intrinsic::x86_avx512_conflict_d_128;
2747 else if (Name[9] ==
'd' && VecWidth == 256)
2748 IID = Intrinsic::x86_avx512_conflict_d_256;
2749 else if (Name[9] ==
'd' && VecWidth == 512)
2750 IID = Intrinsic::x86_avx512_conflict_d_512;
2751 else if (Name[9] ==
'q' && VecWidth == 128)
2752 IID = Intrinsic::x86_avx512_conflict_q_128;
2753 else if (Name[9] ==
'q' && VecWidth == 256)
2754 IID = Intrinsic::x86_avx512_conflict_q_256;
2755 else if (Name[9] ==
'q' && VecWidth == 512)
2756 IID = Intrinsic::x86_avx512_conflict_q_512;
2759 }
else if (Name.starts_with(
"pavg.")) {
2760 if (Name[5] ==
'b' && VecWidth == 128)
2761 IID = Intrinsic::x86_sse2_pavg_b;
2762 else if (Name[5] ==
'b' && VecWidth == 256)
2763 IID = Intrinsic::x86_avx2_pavg_b;
2764 else if (Name[5] ==
'b' && VecWidth == 512)
2765 IID = Intrinsic::x86_avx512_pavg_b_512;
2766 else if (Name[5] ==
'w' && VecWidth == 128)
2767 IID = Intrinsic::x86_sse2_pavg_w;
2768 else if (Name[5] ==
'w' && VecWidth == 256)
2769 IID = Intrinsic::x86_avx2_pavg_w;
2770 else if (Name[5] ==
'w' && VecWidth == 512)
2771 IID = Intrinsic::x86_avx512_pavg_w_512;
2780 Rep = Builder.CreateIntrinsic(IID, Args);
2791 if (AsmStr->find(
"mov\tfp") == 0 &&
2792 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
2793 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
2794 AsmStr->replace(Pos, 1,
";");
2800 Value *Rep =
nullptr;
2802 if (Name ==
"abs.i" || Name ==
"abs.ll") {
2804 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
2805 {Arg, Builder.getTrue()},
2807 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
2808 Type *Ty = (Name ==
"abs.bf16")
2812 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
2813 Rep = Builder.CreateBitCast(Abs, CI->
getType());
2814 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
2815 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
2816 : Intrinsic::nvvm_fabs;
2817 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
2818 }
else if (Name.consume_front(
"ex2.approx.")) {
2820 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
2821 : Intrinsic::nvvm_ex2_approx;
2822 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
2823 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
2824 Name.starts_with(
"atomic.load.add.f64.p")) {
2827 Rep = Builder.CreateAtomicRMW(
2833 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
2834 Name.starts_with(
"atomic.load.dec.32.p")) {
2839 Rep = Builder.CreateAtomicRMW(
2843 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
2849 Op.contains(
".cta.") ?
"block" :
"");
2850 if (
Op.starts_with(
"cas.")) {
2852 Value *Pair = Builder.CreateAtomicCmpXchg(
2855 Rep = Builder.CreateExtractValue(Pair, 0);
2873 "unexpected nvvm scoped atomic intrinsic");
2874 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
2877 }
else if (Name ==
"clz.ll") {
2880 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
2881 {Arg, Builder.getFalse()},
2883 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
2884 }
else if (Name ==
"popc.ll") {
2888 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
2889 Arg,
nullptr,
"ctpop");
2890 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
2891 }
else if (Name ==
"h2f") {
2893 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
2894 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
2895 }
else if (Name.consume_front(
"bitcast.") &&
2896 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
2899 }
else if (Name ==
"rotate.b32") {
2902 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
2903 {Arg, Arg, ShiftAmt});
2904 }
else if (Name ==
"rotate.b64") {
2908 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
2909 {Arg, Arg, ZExtShiftAmt});
2910 }
else if (Name ==
"rotate.right.b64") {
2914 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
2915 {Arg, Arg, ZExtShiftAmt});
2916 }
else if (Name ==
"swap.lo.hi.b64") {
2919 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
2920 {Arg, Arg, Builder.getInt64(32)});
2921 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
2924 Name.starts_with(
".to.gen"))) {
2926 }
else if (Name.consume_front(
"ldg.global")) {
2930 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
2933 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
2935 }
else if (Name ==
"tanh.approx.f32") {
2939 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
2941 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
2943 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
2944 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
2946 }
else if (Name ==
"barrier") {
2947 Rep = Builder.CreateIntrinsic(
2948 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
2950 }
else if (Name ==
"barrier.sync") {
2951 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
2953 }
else if (Name ==
"barrier.sync.cnt") {
2954 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
2956 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
2957 Name ==
"barrier0.or") {
2959 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
2963 .
Case(
"barrier0.popc",
2964 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
2965 .
Case(
"barrier0.and",
2966 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
2967 .
Case(
"barrier0.or",
2968 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
2969 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
2970 Rep = Builder.CreateZExt(Bar, CI->
getType());
2974 !
F->getReturnType()->getScalarType()->isBFloatTy()) {
2984 ? Builder.CreateBitCast(Arg, NewType)
2987 Rep = Builder.CreateCall(NewFn, Args);
2988 if (
F->getReturnType()->isIntegerTy())
2989 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
2999 Value *Rep =
nullptr;
3001 if (Name.starts_with(
"sse4a.movnt.")) {
3013 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3016 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3017 }
else if (Name.starts_with(
"avx.movnt.") ||
3018 Name.starts_with(
"avx512.storent.")) {
3030 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3031 }
else if (Name ==
"sse2.storel.dq") {
3036 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3037 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3038 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3039 }
else if (Name.starts_with(
"sse.storeu.") ||
3040 Name.starts_with(
"sse2.storeu.") ||
3041 Name.starts_with(
"avx.storeu.")) {
3044 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3045 }
else if (Name ==
"avx512.mask.store.ss") {
3049 }
else if (Name.starts_with(
"avx512.mask.store")) {
3051 bool Aligned = Name[17] !=
'u';
3054 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3057 bool CmpEq = Name[9] ==
'e';
3060 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3061 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3068 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3069 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3071 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3072 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3073 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3074 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3075 Name.starts_with(
"sse2.sqrt.p") ||
3076 Name.starts_with(
"sse.sqrt.p")) {
3077 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3078 {CI->getArgOperand(0)});
3079 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3083 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3084 : Intrinsic::x86_avx512_sqrt_pd_512;
3087 Rep = Builder.CreateIntrinsic(IID, Args);
3089 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3090 {CI->getArgOperand(0)});
3094 }
else if (Name.starts_with(
"avx512.ptestm") ||
3095 Name.starts_with(
"avx512.ptestnm")) {
3099 Rep = Builder.CreateAnd(Op0, Op1);
3105 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3107 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3110 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3113 }
else if (Name.starts_with(
"avx512.kunpck")) {
3118 for (
unsigned i = 0; i != NumElts; ++i)
3127 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3128 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3129 }
else if (Name ==
"avx512.kand.w") {
3132 Rep = Builder.CreateAnd(
LHS,
RHS);
3133 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3134 }
else if (Name ==
"avx512.kandn.w") {
3137 LHS = Builder.CreateNot(
LHS);
3138 Rep = Builder.CreateAnd(
LHS,
RHS);
3139 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3140 }
else if (Name ==
"avx512.kor.w") {
3143 Rep = Builder.CreateOr(
LHS,
RHS);
3144 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3145 }
else if (Name ==
"avx512.kxor.w") {
3148 Rep = Builder.CreateXor(
LHS,
RHS);
3149 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3150 }
else if (Name ==
"avx512.kxnor.w") {
3153 LHS = Builder.CreateNot(
LHS);
3154 Rep = Builder.CreateXor(
LHS,
RHS);
3155 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3156 }
else if (Name ==
"avx512.knot.w") {
3158 Rep = Builder.CreateNot(Rep);
3159 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3160 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3163 Rep = Builder.CreateOr(
LHS,
RHS);
3164 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3166 if (Name[14] ==
'c')
3170 Rep = Builder.CreateICmpEQ(Rep,
C);
3171 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3172 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3173 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3174 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3175 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3178 ConstantInt::get(I32Ty, 0));
3180 ConstantInt::get(I32Ty, 0));
3182 if (Name.contains(
".add."))
3183 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3184 else if (Name.contains(
".sub."))
3185 EltOp = Builder.CreateFSub(Elt0, Elt1);
3186 else if (Name.contains(
".mul."))
3187 EltOp = Builder.CreateFMul(Elt0, Elt1);
3189 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3190 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3191 ConstantInt::get(I32Ty, 0));
3192 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3194 bool CmpEq = Name[16] ==
'e';
3196 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3198 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3205 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3208 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3211 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3218 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3220 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3221 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3223 if (VecWidth == 128 && EltWidth == 32)
3224 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3225 else if (VecWidth == 256 && EltWidth == 32)
3226 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3227 else if (VecWidth == 512 && EltWidth == 32)
3228 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3229 else if (VecWidth == 128 && EltWidth == 64)
3230 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3231 else if (VecWidth == 256 && EltWidth == 64)
3232 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3233 else if (VecWidth == 512 && EltWidth == 64)
3234 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3241 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3244 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3245 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3247 if (VecWidth == 128 && EltWidth == 32)
3248 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3249 else if (VecWidth == 256 && EltWidth == 32)
3250 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3251 else if (VecWidth == 512 && EltWidth == 32)
3252 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3253 else if (VecWidth == 128 && EltWidth == 64)
3254 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3255 else if (VecWidth == 256 && EltWidth == 64)
3256 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3257 else if (VecWidth == 512 && EltWidth == 64)
3258 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3263 if (VecWidth == 512)
3265 Args.push_back(Mask);
3267 Rep = Builder.CreateIntrinsic(IID, Args);
3268 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3272 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3275 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3276 Name.starts_with(
"avx512.cvtw2mask.") ||
3277 Name.starts_with(
"avx512.cvtd2mask.") ||
3278 Name.starts_with(
"avx512.cvtq2mask.")) {
3283 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3284 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3285 Name.starts_with(
"avx512.mask.pabs")) {
3287 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3288 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3289 Name.starts_with(
"avx512.mask.pmaxs")) {
3291 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3292 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3293 Name.starts_with(
"avx512.mask.pmaxu")) {
3295 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3296 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3297 Name.starts_with(
"avx512.mask.pmins")) {
3299 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3300 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3301 Name.starts_with(
"avx512.mask.pminu")) {
3303 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3304 Name ==
"avx512.pmulu.dq.512" ||
3305 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3307 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3308 Name ==
"avx512.pmul.dq.512" ||
3309 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3311 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3312 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3317 }
else if (Name ==
"avx512.cvtusi2sd") {
3322 }
else if (Name ==
"sse2.cvtss2sd") {
3324 Rep = Builder.CreateFPExt(
3327 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3328 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3329 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3330 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3331 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3332 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3333 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3334 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3335 Name ==
"avx512.mask.cvtqq2ps.256" ||
3336 Name ==
"avx512.mask.cvtqq2ps.512" ||
3337 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3338 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3339 Name ==
"avx.cvt.ps2.pd.256" ||
3340 Name ==
"avx512.mask.cvtps2pd.128" ||
3341 Name ==
"avx512.mask.cvtps2pd.256") {
3346 unsigned NumDstElts = DstTy->getNumElements();
3347 if (NumDstElts < SrcTy->getNumElements()) {
3348 assert(NumDstElts == 2 &&
"Unexpected vector size");
3349 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3352 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3353 bool IsUnsigned = Name.contains(
"cvtu");
3355 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3359 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3360 : Intrinsic::x86_avx512_sitofp_round;
3361 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3364 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3365 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3371 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3372 Name.starts_with(
"vcvtph2ps.")) {
3376 unsigned NumDstElts = DstTy->getNumElements();
3377 if (NumDstElts != SrcTy->getNumElements()) {
3378 assert(NumDstElts == 4 &&
"Unexpected vector size");
3379 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3381 Rep = Builder.CreateBitCast(
3383 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3387 }
else if (Name.starts_with(
"avx512.mask.load")) {
3389 bool Aligned = Name[16] !=
'u';
3392 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3396 ResultTy->getNumElements());
3397 Rep = Builder.CreateIntrinsic(
3398 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3400 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3406 Rep = Builder.CreateIntrinsic(
3407 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3409 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3410 Name.starts_with(
"avx512.mask.expand.")) {
3414 ResultTy->getNumElements());
3416 bool IsCompress = Name[12] ==
'c';
3417 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3418 : Intrinsic::x86_avx512_mask_expand;
3419 Rep = Builder.CreateIntrinsic(
3421 }
else if (Name.starts_with(
"xop.vpcom")) {
3423 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3424 Name.ends_with(
"uq"))
3426 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3427 Name.ends_with(
"d") || Name.ends_with(
"q"))
3436 Name = Name.substr(9);
3437 if (Name.starts_with(
"lt"))
3439 else if (Name.starts_with(
"le"))
3441 else if (Name.starts_with(
"gt"))
3443 else if (Name.starts_with(
"ge"))
3445 else if (Name.starts_with(
"eq"))
3447 else if (Name.starts_with(
"ne"))
3449 else if (Name.starts_with(
"false"))
3451 else if (Name.starts_with(
"true"))
3458 }
else if (Name.starts_with(
"xop.vpcmov")) {
3460 Value *NotSel = Builder.CreateNot(Sel);
3463 Rep = Builder.CreateOr(Sel0, Sel1);
3464 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3465 Name.starts_with(
"avx512.mask.prol")) {
3467 }
else if (Name.starts_with(
"avx512.pror") ||
3468 Name.starts_with(
"avx512.mask.pror")) {
3470 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3471 Name.starts_with(
"avx512.mask.vpshld") ||
3472 Name.starts_with(
"avx512.maskz.vpshld")) {
3473 bool ZeroMask = Name[11] ==
'z';
3475 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3476 Name.starts_with(
"avx512.mask.vpshrd") ||
3477 Name.starts_with(
"avx512.maskz.vpshrd")) {
3478 bool ZeroMask = Name[11] ==
'z';
3480 }
else if (Name ==
"sse42.crc32.64.8") {
3483 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3485 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3486 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3487 Name.starts_with(
"avx512.vbroadcast.s")) {
3490 Type *EltTy = VecTy->getElementType();
3491 unsigned EltNum = VecTy->getNumElements();
3495 for (
unsigned I = 0;
I < EltNum; ++
I)
3496 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3497 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3498 Name.starts_with(
"sse41.pmovzx") ||
3499 Name.starts_with(
"avx2.pmovsx") ||
3500 Name.starts_with(
"avx2.pmovzx") ||
3501 Name.starts_with(
"avx512.mask.pmovsx") ||
3502 Name.starts_with(
"avx512.mask.pmovzx")) {
3504 unsigned NumDstElts = DstTy->getNumElements();
3508 for (
unsigned i = 0; i != NumDstElts; ++i)
3513 bool DoSext = Name.contains(
"pmovsx");
3515 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3520 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3521 Name ==
"avx512.mask.pmov.qd.512" ||
3522 Name ==
"avx512.mask.pmov.wb.256" ||
3523 Name ==
"avx512.mask.pmov.wb.512") {
3528 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3529 Name ==
"avx2.vbroadcasti128") {
3535 if (NumSrcElts == 2)
3538 Rep = Builder.CreateShuffleVector(
Load,
3540 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3541 Name.starts_with(
"avx512.mask.shuf.f")) {
3546 unsigned ControlBitsMask = NumLanes - 1;
3547 unsigned NumControlBits = NumLanes / 2;
3550 for (
unsigned l = 0; l != NumLanes; ++l) {
3551 unsigned LaneMask = (Imm >> (l * NumControlBits)) & ControlBitsMask;
3553 if (l >= NumLanes / 2)
3554 LaneMask += NumLanes;
3555 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3556 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3562 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3563 Name.starts_with(
"avx512.mask.broadcasti")) {
3566 unsigned NumDstElts =
3570 for (
unsigned i = 0; i != NumDstElts; ++i)
3571 ShuffleMask[i] = i % NumSrcElts;
3577 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3578 Name.starts_with(
"avx2.vbroadcast") ||
3579 Name.starts_with(
"avx512.pbroadcast") ||
3580 Name.starts_with(
"avx512.mask.broadcast.s")) {
3587 Rep = Builder.CreateShuffleVector(
Op, M);
3592 }
else if (Name.starts_with(
"sse2.padds.") ||
3593 Name.starts_with(
"avx2.padds.") ||
3594 Name.starts_with(
"avx512.padds.") ||
3595 Name.starts_with(
"avx512.mask.padds.")) {
3597 }
else if (Name.starts_with(
"sse2.psubs.") ||
3598 Name.starts_with(
"avx2.psubs.") ||
3599 Name.starts_with(
"avx512.psubs.") ||
3600 Name.starts_with(
"avx512.mask.psubs.")) {
3602 }
else if (Name.starts_with(
"sse2.paddus.") ||
3603 Name.starts_with(
"avx2.paddus.") ||
3604 Name.starts_with(
"avx512.mask.paddus.")) {
3606 }
else if (Name.starts_with(
"sse2.psubus.") ||
3607 Name.starts_with(
"avx2.psubus.") ||
3608 Name.starts_with(
"avx512.mask.psubus.")) {
3610 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3615 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3619 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3624 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3629 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3630 Name ==
"avx512.psll.dq.512") {
3634 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3635 Name ==
"avx512.psrl.dq.512") {
3639 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3640 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3641 Name.starts_with(
"avx2.pblendd.")) {
3646 unsigned NumElts = VecTy->getNumElements();
3649 for (
unsigned i = 0; i != NumElts; ++i)
3650 Idxs[i] = ((Imm >> (i % 8)) & 1) ? i + NumElts : i;
3652 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3653 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3654 Name ==
"avx2.vinserti128" ||
3655 Name.starts_with(
"avx512.mask.insert")) {
3659 unsigned DstNumElts =
3661 unsigned SrcNumElts =
3663 unsigned Scale = DstNumElts / SrcNumElts;
3670 for (
unsigned i = 0; i != SrcNumElts; ++i)
3672 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3673 Idxs[i] = SrcNumElts;
3674 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3688 for (
unsigned i = 0; i != DstNumElts; ++i)
3691 for (
unsigned i = 0; i != SrcNumElts; ++i)
3692 Idxs[i + Imm * SrcNumElts] = i + DstNumElts;
3693 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3699 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3700 Name ==
"avx2.vextracti128" ||
3701 Name.starts_with(
"avx512.mask.vextract")) {
3704 unsigned DstNumElts =
3706 unsigned SrcNumElts =
3708 unsigned Scale = SrcNumElts / DstNumElts;
3715 for (
unsigned i = 0; i != DstNumElts; ++i) {
3716 Idxs[i] = i + (Imm * DstNumElts);
3718 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3724 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3725 Name.starts_with(
"avx512.mask.perm.di.")) {
3729 unsigned NumElts = VecTy->getNumElements();
3732 for (
unsigned i = 0; i != NumElts; ++i)
3733 Idxs[i] = (i & ~0x3) + ((Imm >> (2 * (i & 0x3))) & 3);
3735 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3740 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
3752 unsigned HalfSize = NumElts / 2;
3764 unsigned StartIndex = (Imm & 0x01) ? HalfSize : 0;
3765 for (
unsigned i = 0; i < HalfSize; ++i)
3766 ShuffleMask[i] = StartIndex + i;
3769 StartIndex = (Imm & 0x10) ? HalfSize : 0;
3770 for (
unsigned i = 0; i < HalfSize; ++i)
3771 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
3773 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
3775 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
3776 Name.starts_with(
"avx512.mask.vpermil.p") ||
3777 Name.starts_with(
"avx512.mask.pshuf.d.")) {
3781 unsigned NumElts = VecTy->getNumElements();
3783 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
3784 unsigned IdxMask = ((1 << IdxSize) - 1);
3790 for (
unsigned i = 0; i != NumElts; ++i)
3791 Idxs[i] = ((Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
3793 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3798 }
else if (Name ==
"sse2.pshufl.w" ||
3799 Name.starts_with(
"avx512.mask.pshufl.w.")) {
3804 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
3808 for (
unsigned l = 0; l != NumElts; l += 8) {
3809 for (
unsigned i = 0; i != 4; ++i)
3810 Idxs[i + l] = ((Imm >> (2 * i)) & 0x3) + l;
3811 for (
unsigned i = 4; i != 8; ++i)
3812 Idxs[i + l] = i + l;
3815 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3820 }
else if (Name ==
"sse2.pshufh.w" ||
3821 Name.starts_with(
"avx512.mask.pshufh.w.")) {
3826 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
3830 for (
unsigned l = 0; l != NumElts; l += 8) {
3831 for (
unsigned i = 0; i != 4; ++i)
3832 Idxs[i + l] = i + l;
3833 for (
unsigned i = 0; i != 4; ++i)
3834 Idxs[i + l + 4] = ((Imm >> (2 * i)) & 0x3) + 4 + l;
3837 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3842 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
3849 unsigned HalfLaneElts = NumLaneElts / 2;
3852 for (
unsigned i = 0; i != NumElts; ++i) {
3854 Idxs[i] = i - (i % NumLaneElts);
3856 if ((i % NumLaneElts) >= HalfLaneElts)
3860 Idxs[i] += (Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
3863 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3867 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
3868 Name.starts_with(
"avx512.mask.movshdup") ||
3869 Name.starts_with(
"avx512.mask.movsldup")) {
3875 if (Name.starts_with(
"avx512.mask.movshdup."))
3879 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
3880 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
3881 Idxs[i + l + 0] = i + l +
Offset;
3882 Idxs[i + l + 1] = i + l +
Offset;
3885 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3889 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
3890 Name.starts_with(
"avx512.mask.unpckl.")) {
3897 for (
int l = 0; l != NumElts; l += NumLaneElts)
3898 for (
int i = 0; i != NumLaneElts; ++i)
3899 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
3901 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3905 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
3906 Name.starts_with(
"avx512.mask.unpckh.")) {
3913 for (
int l = 0; l != NumElts; l += NumLaneElts)
3914 for (
int i = 0; i != NumLaneElts; ++i)
3915 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
3917 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3921 }
else if (Name.starts_with(
"avx512.mask.and.") ||
3922 Name.starts_with(
"avx512.mask.pand.")) {
3925 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
3927 Rep = Builder.CreateBitCast(Rep, FTy);
3930 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
3931 Name.starts_with(
"avx512.mask.pandn.")) {
3934 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
3935 Rep = Builder.CreateAnd(Rep,
3937 Rep = Builder.CreateBitCast(Rep, FTy);
3940 }
else if (Name.starts_with(
"avx512.mask.or.") ||
3941 Name.starts_with(
"avx512.mask.por.")) {
3944 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
3946 Rep = Builder.CreateBitCast(Rep, FTy);
3949 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
3950 Name.starts_with(
"avx512.mask.pxor.")) {
3953 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
3955 Rep = Builder.CreateBitCast(Rep, FTy);
3958 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
3962 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
3966 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
3970 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
3971 if (Name.ends_with(
".512")) {
3973 if (Name[17] ==
's')
3974 IID = Intrinsic::x86_avx512_add_ps_512;
3976 IID = Intrinsic::x86_avx512_add_pd_512;
3978 Rep = Builder.CreateIntrinsic(
3986 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
3987 if (Name.ends_with(
".512")) {
3989 if (Name[17] ==
's')
3990 IID = Intrinsic::x86_avx512_div_ps_512;
3992 IID = Intrinsic::x86_avx512_div_pd_512;
3994 Rep = Builder.CreateIntrinsic(
4002 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4003 if (Name.ends_with(
".512")) {
4005 if (Name[17] ==
's')
4006 IID = Intrinsic::x86_avx512_mul_ps_512;
4008 IID = Intrinsic::x86_avx512_mul_pd_512;
4010 Rep = Builder.CreateIntrinsic(
4018 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4019 if (Name.ends_with(
".512")) {
4021 if (Name[17] ==
's')
4022 IID = Intrinsic::x86_avx512_sub_ps_512;
4024 IID = Intrinsic::x86_avx512_sub_pd_512;
4026 Rep = Builder.CreateIntrinsic(
4034 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4035 Name.starts_with(
"avx512.mask.min.p")) &&
4036 Name.drop_front(18) ==
".512") {
4037 bool IsDouble = Name[17] ==
'd';
4038 bool IsMin = Name[13] ==
'i';
4040 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4041 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4044 Rep = Builder.CreateIntrinsic(
4049 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4051 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4052 {CI->getArgOperand(0), Builder.getInt1(false)});
4055 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4056 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4057 bool IsVariable = Name[16] ==
'v';
4058 char Size = Name[16] ==
'.' ? Name[17]
4059 : Name[17] ==
'.' ? Name[18]
4060 : Name[18] ==
'.' ? Name[19]
4064 if (IsVariable && Name[17] !=
'.') {
4065 if (
Size ==
'd' && Name[17] ==
'2')
4066 IID = Intrinsic::x86_avx2_psllv_q;
4067 else if (
Size ==
'd' && Name[17] ==
'4')
4068 IID = Intrinsic::x86_avx2_psllv_q_256;
4069 else if (
Size ==
's' && Name[17] ==
'4')
4070 IID = Intrinsic::x86_avx2_psllv_d;
4071 else if (
Size ==
's' && Name[17] ==
'8')
4072 IID = Intrinsic::x86_avx2_psllv_d_256;
4073 else if (
Size ==
'h' && Name[17] ==
'8')
4074 IID = Intrinsic::x86_avx512_psllv_w_128;
4075 else if (
Size ==
'h' && Name[17] ==
'1')
4076 IID = Intrinsic::x86_avx512_psllv_w_256;
4077 else if (Name[17] ==
'3' && Name[18] ==
'2')
4078 IID = Intrinsic::x86_avx512_psllv_w_512;
4081 }
else if (Name.ends_with(
".128")) {
4083 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4084 : Intrinsic::x86_sse2_psll_d;
4085 else if (
Size ==
'q')
4086 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4087 : Intrinsic::x86_sse2_psll_q;
4088 else if (
Size ==
'w')
4089 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4090 : Intrinsic::x86_sse2_psll_w;
4093 }
else if (Name.ends_with(
".256")) {
4095 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4096 : Intrinsic::x86_avx2_psll_d;
4097 else if (
Size ==
'q')
4098 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4099 : Intrinsic::x86_avx2_psll_q;
4100 else if (
Size ==
'w')
4101 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4102 : Intrinsic::x86_avx2_psll_w;
4107 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4108 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4109 : Intrinsic::x86_avx512_psll_d_512;
4110 else if (
Size ==
'q')
4111 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4112 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4113 : Intrinsic::x86_avx512_psll_q_512;
4114 else if (
Size ==
'w')
4115 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4116 : Intrinsic::x86_avx512_psll_w_512;
4122 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4123 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4124 bool IsVariable = Name[16] ==
'v';
4125 char Size = Name[16] ==
'.' ? Name[17]
4126 : Name[17] ==
'.' ? Name[18]
4127 : Name[18] ==
'.' ? Name[19]
4131 if (IsVariable && Name[17] !=
'.') {
4132 if (
Size ==
'd' && Name[17] ==
'2')
4133 IID = Intrinsic::x86_avx2_psrlv_q;
4134 else if (
Size ==
'd' && Name[17] ==
'4')
4135 IID = Intrinsic::x86_avx2_psrlv_q_256;
4136 else if (
Size ==
's' && Name[17] ==
'4')
4137 IID = Intrinsic::x86_avx2_psrlv_d;
4138 else if (
Size ==
's' && Name[17] ==
'8')
4139 IID = Intrinsic::x86_avx2_psrlv_d_256;
4140 else if (
Size ==
'h' && Name[17] ==
'8')
4141 IID = Intrinsic::x86_avx512_psrlv_w_128;
4142 else if (
Size ==
'h' && Name[17] ==
'1')
4143 IID = Intrinsic::x86_avx512_psrlv_w_256;
4144 else if (Name[17] ==
'3' && Name[18] ==
'2')
4145 IID = Intrinsic::x86_avx512_psrlv_w_512;
4148 }
else if (Name.ends_with(
".128")) {
4150 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4151 : Intrinsic::x86_sse2_psrl_d;
4152 else if (
Size ==
'q')
4153 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4154 : Intrinsic::x86_sse2_psrl_q;
4155 else if (
Size ==
'w')
4156 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4157 : Intrinsic::x86_sse2_psrl_w;
4160 }
else if (Name.ends_with(
".256")) {
4162 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4163 : Intrinsic::x86_avx2_psrl_d;
4164 else if (
Size ==
'q')
4165 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4166 : Intrinsic::x86_avx2_psrl_q;
4167 else if (
Size ==
'w')
4168 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4169 : Intrinsic::x86_avx2_psrl_w;
4174 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4175 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4176 : Intrinsic::x86_avx512_psrl_d_512;
4177 else if (
Size ==
'q')
4178 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4179 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4180 : Intrinsic::x86_avx512_psrl_q_512;
4181 else if (
Size ==
'w')
4182 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4183 : Intrinsic::x86_avx512_psrl_w_512;
4189 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4190 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4191 bool IsVariable = Name[16] ==
'v';
4192 char Size = Name[16] ==
'.' ? Name[17]
4193 : Name[17] ==
'.' ? Name[18]
4194 : Name[18] ==
'.' ? Name[19]
4198 if (IsVariable && Name[17] !=
'.') {
4199 if (
Size ==
's' && Name[17] ==
'4')
4200 IID = Intrinsic::x86_avx2_psrav_d;
4201 else if (
Size ==
's' && Name[17] ==
'8')
4202 IID = Intrinsic::x86_avx2_psrav_d_256;
4203 else if (
Size ==
'h' && Name[17] ==
'8')
4204 IID = Intrinsic::x86_avx512_psrav_w_128;
4205 else if (
Size ==
'h' && Name[17] ==
'1')
4206 IID = Intrinsic::x86_avx512_psrav_w_256;
4207 else if (Name[17] ==
'3' && Name[18] ==
'2')
4208 IID = Intrinsic::x86_avx512_psrav_w_512;
4211 }
else if (Name.ends_with(
".128")) {
4213 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4214 : Intrinsic::x86_sse2_psra_d;
4215 else if (
Size ==
'q')
4216 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4217 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4218 : Intrinsic::x86_avx512_psra_q_128;
4219 else if (
Size ==
'w')
4220 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4221 : Intrinsic::x86_sse2_psra_w;
4224 }
else if (Name.ends_with(
".256")) {
4226 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4227 : Intrinsic::x86_avx2_psra_d;
4228 else if (
Size ==
'q')
4229 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4230 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4231 : Intrinsic::x86_avx512_psra_q_256;
4232 else if (
Size ==
'w')
4233 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4234 : Intrinsic::x86_avx2_psra_w;
4239 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4240 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4241 : Intrinsic::x86_avx512_psra_d_512;
4242 else if (
Size ==
'q')
4243 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4244 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4245 : Intrinsic::x86_avx512_psra_q_512;
4246 else if (
Size ==
'w')
4247 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4248 : Intrinsic::x86_avx512_psra_w_512;
4254 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4256 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4258 }
else if (Name.ends_with(
".movntdqa")) {
4262 LoadInst *LI = Builder.CreateAlignedLoad(
4267 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4268 Name.starts_with(
"fma.vfmsub.") ||
4269 Name.starts_with(
"fma.vfnmadd.") ||
4270 Name.starts_with(
"fma.vfnmsub.")) {
4271 bool NegMul = Name[6] ==
'n';
4272 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4273 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4284 if (NegMul && !IsScalar)
4285 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4286 if (NegMul && IsScalar)
4287 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4289 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4291 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4295 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4303 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4307 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4308 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4309 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4310 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4311 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4312 bool IsMask3 = Name[11] ==
'3';
4313 bool IsMaskZ = Name[11] ==
'z';
4315 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4316 bool NegMul = Name[2] ==
'n';
4317 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4323 if (NegMul && (IsMask3 || IsMaskZ))
4324 A = Builder.CreateFNeg(
A);
4325 if (NegMul && !(IsMask3 || IsMaskZ))
4326 B = Builder.CreateFNeg(
B);
4328 C = Builder.CreateFNeg(
C);
4330 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4331 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4332 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4339 if (Name.back() ==
'd')
4340 IID = Intrinsic::x86_avx512_vfmadd_f64;
4342 IID = Intrinsic::x86_avx512_vfmadd_f32;
4343 Rep = Builder.CreateIntrinsic(IID,
Ops);
4345 Rep = Builder.CreateFMA(
A,
B,
C);
4354 if (NegAcc && IsMask3)
4359 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4361 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4362 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4363 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4364 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4365 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4366 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4367 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4368 bool IsMask3 = Name[11] ==
'3';
4369 bool IsMaskZ = Name[11] ==
'z';
4371 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4372 bool NegMul = Name[2] ==
'n';
4373 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4379 if (NegMul && (IsMask3 || IsMaskZ))
4380 A = Builder.CreateFNeg(
A);
4381 if (NegMul && !(IsMask3 || IsMaskZ))
4382 B = Builder.CreateFNeg(
B);
4384 C = Builder.CreateFNeg(
C);
4391 if (Name[Name.size() - 5] ==
's')
4392 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4394 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4398 Rep = Builder.CreateFMA(
A,
B,
C);
4406 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4410 if (VecWidth == 128 && EltWidth == 32)
4411 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4412 else if (VecWidth == 256 && EltWidth == 32)
4413 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4414 else if (VecWidth == 128 && EltWidth == 64)
4415 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4416 else if (VecWidth == 256 && EltWidth == 64)
4417 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4423 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4424 Rep = Builder.CreateIntrinsic(IID,
Ops);
4425 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4426 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4427 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4428 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4429 bool IsMask3 = Name[11] ==
'3';
4430 bool IsMaskZ = Name[11] ==
'z';
4432 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4433 bool IsSubAdd = Name[3] ==
's';
4437 if (Name[Name.size() - 5] ==
's')
4438 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4440 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4445 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4447 Rep = Builder.CreateIntrinsic(IID,
Ops);
4456 Value *Odd = Builder.CreateCall(FMA,
Ops);
4457 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4458 Value *Even = Builder.CreateCall(FMA,
Ops);
4464 for (
int i = 0; i != NumElts; ++i)
4465 Idxs[i] = i + (i % 2) * NumElts;
4467 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4475 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4476 Name.starts_with(
"avx512.maskz.pternlog.")) {
4477 bool ZeroMask = Name[11] ==
'z';
4481 if (VecWidth == 128 && EltWidth == 32)
4482 IID = Intrinsic::x86_avx512_pternlog_d_128;
4483 else if (VecWidth == 256 && EltWidth == 32)
4484 IID = Intrinsic::x86_avx512_pternlog_d_256;
4485 else if (VecWidth == 512 && EltWidth == 32)
4486 IID = Intrinsic::x86_avx512_pternlog_d_512;
4487 else if (VecWidth == 128 && EltWidth == 64)
4488 IID = Intrinsic::x86_avx512_pternlog_q_128;
4489 else if (VecWidth == 256 && EltWidth == 64)
4490 IID = Intrinsic::x86_avx512_pternlog_q_256;
4491 else if (VecWidth == 512 && EltWidth == 64)
4492 IID = Intrinsic::x86_avx512_pternlog_q_512;
4498 Rep = Builder.CreateIntrinsic(IID, Args);
4502 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4503 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4504 bool ZeroMask = Name[11] ==
'z';
4505 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4508 if (VecWidth == 128 && !
High)
4509 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4510 else if (VecWidth == 256 && !
High)
4511 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4512 else if (VecWidth == 512 && !
High)
4513 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4514 else if (VecWidth == 128 &&
High)
4515 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4516 else if (VecWidth == 256 &&
High)
4517 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4518 else if (VecWidth == 512 &&
High)
4519 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4525 Rep = Builder.CreateIntrinsic(IID, Args);
4529 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4530 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4531 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4532 bool ZeroMask = Name[11] ==
'z';
4533 bool IndexForm = Name[17] ==
'i';
4535 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4536 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4537 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4538 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4539 bool ZeroMask = Name[11] ==
'z';
4540 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4543 if (VecWidth == 128 && !IsSaturating)
4544 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4545 else if (VecWidth == 256 && !IsSaturating)
4546 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4547 else if (VecWidth == 512 && !IsSaturating)
4548 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4549 else if (VecWidth == 128 && IsSaturating)
4550 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4551 else if (VecWidth == 256 && IsSaturating)
4552 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4553 else if (VecWidth == 512 && IsSaturating)
4554 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4564 if (Args[1]->
getType()->isVectorTy() &&
4567 ->isIntegerTy(32) &&
4568 Args[2]->
getType()->isVectorTy() &&
4571 ->isIntegerTy(32)) {
4572 Type *NewArgType =
nullptr;
4573 if (VecWidth == 128)
4575 else if (VecWidth == 256)
4577 else if (VecWidth == 512)
4583 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4584 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4587 Rep = Builder.CreateIntrinsic(IID, Args);
4591 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4592 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4593 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4594 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4595 bool ZeroMask = Name[11] ==
'z';
4596 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4599 if (VecWidth == 128 && !IsSaturating)
4600 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4601 else if (VecWidth == 256 && !IsSaturating)
4602 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4603 else if (VecWidth == 512 && !IsSaturating)
4604 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4605 else if (VecWidth == 128 && IsSaturating)
4606 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4607 else if (VecWidth == 256 && IsSaturating)
4608 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4609 else if (VecWidth == 512 && IsSaturating)
4610 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4620 if (Args[1]->
getType()->isVectorTy() &&
4623 ->isIntegerTy(32) &&
4624 Args[2]->
getType()->isVectorTy() &&
4627 ->isIntegerTy(32)) {
4628 Type *NewArgType =
nullptr;
4629 if (VecWidth == 128)
4631 else if (VecWidth == 256)
4633 else if (VecWidth == 512)
4639 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4640 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4643 Rep = Builder.CreateIntrinsic(IID, Args);
4647 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4648 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4649 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4651 if (Name[0] ==
'a' && Name.back() ==
'2')
4652 IID = Intrinsic::x86_addcarry_32;
4653 else if (Name[0] ==
'a' && Name.back() ==
'4')
4654 IID = Intrinsic::x86_addcarry_64;
4655 else if (Name[0] ==
's' && Name.back() ==
'2')
4656 IID = Intrinsic::x86_subborrow_32;
4657 else if (Name[0] ==
's' && Name.back() ==
'4')
4658 IID = Intrinsic::x86_subborrow_64;
4665 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4668 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4671 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4675 }
else if (Name.starts_with(
"avx512.mask.") &&
4678 }
else if (Name.starts_with(
"bmi.pdep.")) {
4680 }
else if (Name.starts_with(
"bmi.pext.")) {
4690 if (Name.starts_with(
"neon.bfcvt")) {
4691 if (Name.starts_with(
"neon.bfcvtn2")) {
4693 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4695 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4696 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4699 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4700 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4702 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4706 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4707 return Builder.CreateShuffleVector(
4710 return Builder.CreateFPTrunc(CI->
getOperand(0),
4713 }
else if (Name.starts_with(
"sve.fcvt")) {
4716 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4717 .
Case(
"sve.fcvtnt.bf16f32",
4718 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4730 if (Args[1]->
getType() != BadPredTy)
4733 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
4734 BadPredTy, Args[1]);
4735 Args[1] = Builder.CreateIntrinsic(
4736 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
4738 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
4742 if (Name ==
"neon.vcvtfp2hf")
4743 return Builder.CreateBitCast(
4744 Builder.CreateFPTrunc(
4748 if (Name ==
"neon.vcvthf2fp")
4749 return Builder.CreateFPExt(
4750 Builder.CreateBitCast(
4760 if (Name ==
"mve.vctp64.old") {
4763 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
4766 Value *C1 = Builder.CreateIntrinsic(
4767 Intrinsic::arm_mve_pred_v2i,
4769 return Builder.CreateIntrinsic(
4770 Intrinsic::arm_mve_pred_i2v,
4772 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
4773 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
4774 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
4775 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
4777 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
4778 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
4779 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
4780 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
4782 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
4783 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
4784 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
4785 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
4786 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
4787 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
4788 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
4789 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
4790 std::vector<Type *> Tys;
4794 case Intrinsic::arm_mve_mull_int_predicated:
4795 case Intrinsic::arm_mve_vqdmull_predicated:
4796 case Intrinsic::arm_mve_vldr_gather_base_predicated:
4799 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
4800 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
4801 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
4805 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
4809 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
4813 case Intrinsic::arm_cde_vcx1q_predicated:
4814 case Intrinsic::arm_cde_vcx1qa_predicated:
4815 case Intrinsic::arm_cde_vcx2q_predicated:
4816 case Intrinsic::arm_cde_vcx2qa_predicated:
4817 case Intrinsic::arm_cde_vcx3q_predicated:
4818 case Intrinsic::arm_cde_vcx3qa_predicated:
4825 std::vector<Value *>
Ops;
4827 Type *Ty =
Op->getType();
4828 if (Ty->getScalarSizeInBits() == 1) {
4829 Value *C1 = Builder.CreateIntrinsic(
4830 Intrinsic::arm_mve_pred_v2i,
4832 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
4837 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
4852 auto UpgradeLegacyWMMAIUIntrinsicCall =
4857 Args.push_back(Builder.getFalse());
4861 F->getParent(),
F->getIntrinsicID(), OverloadTys);
4868 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
4873 NewCall->copyMetadata(*CI);
4877 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
4878 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
4879 "intrinsic should have 7 arguments");
4882 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
4884 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
4885 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
4886 "intrinsic should have 8 arguments");
4891 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
4894 switch (
F->getIntrinsicID()) {
4897 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
4898 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
4899 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
4900 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
4901 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
4902 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
4917 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
4920 F->getParent(),
F->getIntrinsicID(), Overloads);
4925 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
4930 NewCall->copyMetadata(*CI);
4931 NewCall->takeName(CI);
4953 if (NumOperands < 3)
4966 bool IsVolatile =
false;
4970 if (NumOperands > 3)
4975 if (NumOperands > 5) {
4977 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
4991 if (VT->getElementType()->isIntegerTy(16)) {
4994 Val = Builder.CreateBitCast(Val, AsBF16);
5002 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5004 unsigned AddrSpace = PtrTy->getAddressSpace();
5007 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5009 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5014 MDNode *RangeNotPrivate =
5017 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5023 return Builder.CreateBitCast(RMW, RetTy);
5044 return MAV->getMetadata();
5053 if (Name ==
"label") {
5055 }
else if (Name ==
"assign") {
5062 }
else if (Name ==
"declare") {
5066 }
else if (Name ==
"addr") {
5076 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5077 }
else if (Name ==
"value") {
5080 unsigned ExprOp = 2;
5095 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5103 int64_t OffsetVal =
Offset->getSExtValue();
5104 return Builder.CreateIntrinsic(OffsetVal >= 0
5105 ? Intrinsic::vector_splice_left
5106 : Intrinsic::vector_splice_right,
5108 {CI->getArgOperand(0), CI->getArgOperand(1),
5109 Builder.getInt32(std::abs(OffsetVal))});
5114 if (Name.starts_with(
"to.fp16")) {
5116 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5117 return Builder.CreateBitCast(Cast, CI->
getType());
5120 if (Name.starts_with(
"from.fp16")) {
5122 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5123 return Builder.CreateFPExt(Cast, CI->
getType());
5134 if (Defaults.empty())
5137 unsigned OldArgCount = CI->
arg_size();
5138 unsigned NewArgCount = NewFn->
arg_size();
5142 if (OldArgCount >= NewArgCount)
5150 if (OldArgCount < FirstDefault)
5155 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5156 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5157 "missing argument outside the default range");
5158 Type *ParamTy = NewFT->getParamType(Idx);
5163 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5169 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5201 if (!Name.consume_front(
"llvm."))
5204 bool IsX86 = Name.consume_front(
"x86.");
5205 bool IsNVVM = Name.consume_front(
"nvvm.");
5206 bool IsAArch64 = Name.consume_front(
"aarch64.");
5207 bool IsARM = Name.consume_front(
"arm.");
5208 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5209 bool IsDbg = Name.consume_front(
"dbg.");
5211 (Name.consume_front(
"experimental.vector.splice") ||
5212 Name.consume_front(
"vector.splice")) &&
5213 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5214 Value *Rep =
nullptr;
5216 if (!IsX86 && Name ==
"stackprotectorcheck") {
5218 }
else if (IsNVVM) {
5222 }
else if (IsAArch64) {
5226 }
else if (IsAMDGCN) {
5230 }
else if (IsOldSplice) {
5232 }
else if (Name.consume_front(
"convert.")) {
5234 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5247 const auto &DefaultCase = [&]() ->
void {
5255 "Unknown function for CallBase upgrade and isn't just a name change");
5263 "Return type must have changed");
5264 assert(OldST->getNumElements() ==
5266 "Must have same number of elements");
5269 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5272 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5273 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5274 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5298 case Intrinsic::arm_neon_vst1:
5299 case Intrinsic::arm_neon_vst2:
5300 case Intrinsic::arm_neon_vst3:
5301 case Intrinsic::arm_neon_vst4:
5302 case Intrinsic::arm_neon_vst2lane:
5303 case Intrinsic::arm_neon_vst3lane:
5304 case Intrinsic::arm_neon_vst4lane: {
5306 NewCall = Builder.CreateCall(NewFn, Args);
5309 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5310 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5311 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5316 NewCall = Builder.CreateCall(NewFn, Args);
5319 case Intrinsic::aarch64_sve_ld3_sret:
5320 case Intrinsic::aarch64_sve_ld4_sret:
5321 case Intrinsic::aarch64_sve_ld2_sret: {
5329 Name = Name.substr(5);
5336 unsigned MinElts = RetTy->getMinNumElements() /
N;
5338 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5340 for (
unsigned I = 0;
I <
N;
I++) {
5341 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5342 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5348 case Intrinsic::coro_end_async:
5349 case Intrinsic::coro_end: {
5351 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5353 NewCall = Builder.CreateCall(NewFn, Args);
5358 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5359 Value *InRamp = Builder.CreateCall(IsInRamp);
5369 case Intrinsic::vector_extract: {
5371 Name = Name.substr(5);
5372 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5377 unsigned MinElts = RetTy->getMinNumElements();
5380 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5384 case Intrinsic::vector_insert: {
5386 Name = Name.substr(5);
5387 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5391 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5396 NewCall = Builder.CreateCall(
5400 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5406 assert(
N > 1 &&
"Create is expected to be between 2-4");
5409 unsigned MinElts = RetTy->getMinNumElements() /
N;
5410 for (
unsigned I = 0;
I <
N;
I++) {
5412 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5419 case Intrinsic::arm_neon_bfdot:
5420 case Intrinsic::arm_neon_bfmmla:
5421 case Intrinsic::arm_neon_bfmlalb:
5422 case Intrinsic::arm_neon_bfmlalt:
5423 case Intrinsic::aarch64_neon_bfdot:
5424 case Intrinsic::aarch64_neon_bfmmla:
5425 case Intrinsic::aarch64_neon_bfmlalb:
5426 case Intrinsic::aarch64_neon_bfmlalt: {
5429 "Mismatch between function args and call args");
5430 size_t OperandWidth =
5432 assert((OperandWidth == 64 || OperandWidth == 128) &&
5433 "Unexpected operand width");
5435 auto Iter = CI->
args().begin();
5436 Args.push_back(*Iter++);
5437 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5438 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5439 NewCall = Builder.CreateCall(NewFn, Args);
5443 case Intrinsic::bitreverse:
5444 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5447 case Intrinsic::ctlz:
5448 case Intrinsic::cttz: {
5455 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5459 case Intrinsic::objectsize: {
5460 Value *NullIsUnknownSize =
5464 NewCall = Builder.CreateCall(
5469 case Intrinsic::ctpop:
5470 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5472 case Intrinsic::dbg_value: {
5474 Name = Name.substr(5);
5476 if (Name.starts_with(
"dbg.addr")) {
5490 if (
Offset->isNullValue()) {
5491 NewCall = Builder.CreateCall(
5500 case Intrinsic::ptr_annotation:
5508 NewCall = Builder.CreateCall(
5517 case Intrinsic::var_annotation:
5524 NewCall = Builder.CreateCall(
5533 case Intrinsic::riscv_aes32dsi:
5534 case Intrinsic::riscv_aes32dsmi:
5535 case Intrinsic::riscv_aes32esi:
5536 case Intrinsic::riscv_aes32esmi:
5537 case Intrinsic::riscv_sm4ks:
5538 case Intrinsic::riscv_sm4ed: {
5548 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5549 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5555 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5556 Value *Res = NewCall;
5558 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5564 case Intrinsic::nvvm_mapa_shared_cluster: {
5568 Value *Res = NewCall;
5569 Res = Builder.CreateAddrSpaceCast(
5576 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5577 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5580 Args[0] = Builder.CreateAddrSpaceCast(
5583 NewCall = Builder.CreateCall(NewFn, Args);
5589 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5590 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5591 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5592 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5593 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5594 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5595 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5596 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5603 Args[0] = Builder.CreateAddrSpaceCast(
5612 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5614 NewCall = Builder.CreateCall(NewFn, Args);
5620 case Intrinsic::riscv_sha256sig0:
5621 case Intrinsic::riscv_sha256sig1:
5622 case Intrinsic::riscv_sha256sum0:
5623 case Intrinsic::riscv_sha256sum1:
5624 case Intrinsic::riscv_sm3p0:
5625 case Intrinsic::riscv_sm3p1: {
5632 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
5634 NewCall = Builder.CreateCall(NewFn, Arg);
5636 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5643 case Intrinsic::x86_xop_vfrcz_ss:
5644 case Intrinsic::x86_xop_vfrcz_sd:
5645 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
5648 case Intrinsic::x86_xop_vpermil2pd:
5649 case Intrinsic::x86_xop_vpermil2ps:
5650 case Intrinsic::x86_xop_vpermil2pd_256:
5651 case Intrinsic::x86_xop_vpermil2ps_256: {
5655 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
5656 NewCall = Builder.CreateCall(NewFn, Args);
5660 case Intrinsic::x86_sse41_ptestc:
5661 case Intrinsic::x86_sse41_ptestz:
5662 case Intrinsic::x86_sse41_ptestnzc: {
5676 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
5677 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
5679 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
5683 case Intrinsic::x86_rdtscp: {
5689 NewCall = Builder.CreateCall(NewFn);
5691 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5694 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
5702 case Intrinsic::x86_sse41_insertps:
5703 case Intrinsic::x86_sse41_dppd:
5704 case Intrinsic::x86_sse41_dpps:
5705 case Intrinsic::x86_sse41_mpsadbw:
5706 case Intrinsic::x86_avx_dp_ps_256:
5707 case Intrinsic::x86_avx2_mpsadbw: {
5713 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
5714 NewCall = Builder.CreateCall(NewFn, Args);
5718 case Intrinsic::x86_avx512_mask_cmp_pd_128:
5719 case Intrinsic::x86_avx512_mask_cmp_pd_256:
5720 case Intrinsic::x86_avx512_mask_cmp_pd_512:
5721 case Intrinsic::x86_avx512_mask_cmp_ps_128:
5722 case Intrinsic::x86_avx512_mask_cmp_ps_256:
5723 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
5729 NewCall = Builder.CreateCall(NewFn, Args);
5738 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
5739 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
5740 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
5741 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
5742 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
5743 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
5747 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
5748 Args[1] = Builder.CreateBitCast(
5751 NewCall = Builder.CreateCall(NewFn, Args);
5752 Value *Res = Builder.CreateBitCast(
5760 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
5761 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
5762 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
5766 Args[1] = Builder.CreateBitCast(
5768 Args[2] = Builder.CreateBitCast(
5771 NewCall = Builder.CreateCall(NewFn, Args);
5775 case Intrinsic::thread_pointer: {
5776 NewCall = Builder.CreateCall(NewFn, {});
5780 case Intrinsic::memcpy:
5781 case Intrinsic::memmove:
5782 case Intrinsic::memset: {
5798 NewCall = Builder.CreateCall(NewFn, Args);
5800 AttributeList NewAttrs = AttributeList::get(
5801 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
5802 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
5803 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
5808 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
5811 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
5815 case Intrinsic::masked_load:
5816 case Intrinsic::masked_gather:
5817 case Intrinsic::masked_store:
5818 case Intrinsic::masked_scatter: {
5824 auto GetMaybeAlign = [](
Value *
Op) {
5834 auto GetAlign = [&](
Value *
Op) {
5843 case Intrinsic::masked_load:
5844 NewCall = Builder.CreateMaskedLoad(
5848 case Intrinsic::masked_gather:
5849 NewCall = Builder.CreateMaskedGather(
5855 case Intrinsic::masked_store:
5856 NewCall = Builder.CreateMaskedStore(
5860 case Intrinsic::masked_scatter:
5861 NewCall = Builder.CreateMaskedScatter(
5863 DL.getValueOrABITypeAlignment(
5877 case Intrinsic::lifetime_start:
5878 case Intrinsic::lifetime_end: {
5890 NewCall = Builder.CreateLifetimeStart(Ptr);
5892 NewCall = Builder.CreateLifetimeEnd(Ptr);
5901 case Intrinsic::x86_avx512_vpdpbusd_128:
5902 case Intrinsic::x86_avx512_vpdpbusd_256:
5903 case Intrinsic::x86_avx512_vpdpbusd_512:
5904 case Intrinsic::x86_avx512_vpdpbusds_128:
5905 case Intrinsic::x86_avx512_vpdpbusds_256:
5906 case Intrinsic::x86_avx512_vpdpbusds_512:
5907 case Intrinsic::x86_avx2_vpdpbssd_128:
5908 case Intrinsic::x86_avx2_vpdpbssd_256:
5909 case Intrinsic::x86_avx10_vpdpbssd_512:
5910 case Intrinsic::x86_avx2_vpdpbssds_128:
5911 case Intrinsic::x86_avx2_vpdpbssds_256:
5912 case Intrinsic::x86_avx10_vpdpbssds_512:
5913 case Intrinsic::x86_avx2_vpdpbsud_128:
5914 case Intrinsic::x86_avx2_vpdpbsud_256:
5915 case Intrinsic::x86_avx10_vpdpbsud_512:
5916 case Intrinsic::x86_avx2_vpdpbsuds_128:
5917 case Intrinsic::x86_avx2_vpdpbsuds_256:
5918 case Intrinsic::x86_avx10_vpdpbsuds_512:
5919 case Intrinsic::x86_avx2_vpdpbuud_128:
5920 case Intrinsic::x86_avx2_vpdpbuud_256:
5921 case Intrinsic::x86_avx10_vpdpbuud_512:
5922 case Intrinsic::x86_avx2_vpdpbuuds_128:
5923 case Intrinsic::x86_avx2_vpdpbuuds_256:
5924 case Intrinsic::x86_avx10_vpdpbuuds_512: {
5929 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5930 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5932 NewCall = Builder.CreateCall(NewFn, Args);
5935 case Intrinsic::x86_avx512_vpdpwssd_128:
5936 case Intrinsic::x86_avx512_vpdpwssd_256:
5937 case Intrinsic::x86_avx512_vpdpwssd_512:
5938 case Intrinsic::x86_avx512_vpdpwssds_128:
5939 case Intrinsic::x86_avx512_vpdpwssds_256:
5940 case Intrinsic::x86_avx512_vpdpwssds_512:
5941 case Intrinsic::x86_avx2_vpdpwsud_128:
5942 case Intrinsic::x86_avx2_vpdpwsud_256:
5943 case Intrinsic::x86_avx10_vpdpwsud_512:
5944 case Intrinsic::x86_avx2_vpdpwsuds_128:
5945 case Intrinsic::x86_avx2_vpdpwsuds_256:
5946 case Intrinsic::x86_avx10_vpdpwsuds_512:
5947 case Intrinsic::x86_avx2_vpdpwusd_128:
5948 case Intrinsic::x86_avx2_vpdpwusd_256:
5949 case Intrinsic::x86_avx10_vpdpwusd_512:
5950 case Intrinsic::x86_avx2_vpdpwusds_128:
5951 case Intrinsic::x86_avx2_vpdpwusds_256:
5952 case Intrinsic::x86_avx10_vpdpwusds_512:
5953 case Intrinsic::x86_avx2_vpdpwuud_128:
5954 case Intrinsic::x86_avx2_vpdpwuud_256:
5955 case Intrinsic::x86_avx10_vpdpwuud_512:
5956 case Intrinsic::x86_avx2_vpdpwuuds_128:
5957 case Intrinsic::x86_avx2_vpdpwuuds_256:
5958 case Intrinsic::x86_avx10_vpdpwuuds_512:
5963 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5964 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5966 NewCall = Builder.CreateCall(NewFn, Args);
5969 assert(NewCall &&
"Should have either set this variable or returned through "
5970 "the default case");
5977 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
5991 F->eraseFromParent();
5997 if (NumOperands == 0)
6005 if (NumOperands == 3) {
6009 Metadata *Elts2[] = {ScalarType, ScalarType,
6023 if (
Opc != Instruction::BitCast)
6027 Type *SrcTy = V->getType();
6044 if (
Opc != Instruction::BitCast)
6047 Type *SrcTy =
C->getType();
6064 if (Flag.getNumOperands() < 3)
6065 return std::nullopt;
6067 return Name->getString();
6068 return std::nullopt;
6082 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6083 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6084 if (auto Name = getModuleFlagNameSafely(*Flag))
6085 return *Name ==
"Debug Info Version";
6088 if (OpIt != ModFlags->op_end()) {
6089 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6096 bool BrokenDebugInfo =
false;
6099 if (!BrokenDebugInfo)
6105 M.getContext().diagnose(Diag);
6112 M.getContext().diagnose(DiagVersion);
6122 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6125 if (
F->hasFnAttribute(Attr)) {
6128 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6130 auto [Part, Rest] = S.
split(
',');
6136 const unsigned Dim = DimC -
'x';
6137 assert(Dim < 3 &&
"Unexpected dim char");
6147 F->addFnAttr(Attr, NewAttr);
6151 return S ==
"x" || S ==
"y" || S ==
"z";
6156 if (K ==
"kernel") {
6168 const unsigned Idx = (AlignIdxValuePair >> 16);
6169 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6174 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6179 if (K ==
"minctasm") {
6184 if (K ==
"maxnreg") {
6189 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6193 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6197 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6201 if (K ==
"grid_constant") {
6216 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6223 if (!SeenNodes.
insert(MD).second)
6230 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6237 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6239 const MDOperand &V = MD->getOperand(j + 1);
6242 NewOperands.
append({K, V});
6245 if (NewOperands.
size() > 1)
6258 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6259 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6260 if (ModRetainReleaseMarker) {
6266 ID->getString().split(ValueComp,
"#");
6267 if (ValueComp.
size() == 2) {
6268 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6272 M.eraseNamedMetadata(ModRetainReleaseMarker);
6283 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6309 bool InvalidCast =
false;
6311 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6324 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6326 Args.push_back(Arg);
6333 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6338 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6351 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6359 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6360 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6361 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6362 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6363 {
"objc_autoreleaseReturnValue",
6364 llvm::Intrinsic::objc_autoreleaseReturnValue},
6365 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6366 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6367 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6368 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6369 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6370 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6371 {
"objc_release", llvm::Intrinsic::objc_release},
6372 {
"objc_retain", llvm::Intrinsic::objc_retain},
6373 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6374 {
"objc_retainAutoreleaseReturnValue",
6375 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6376 {
"objc_retainAutoreleasedReturnValue",
6377 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6378 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6379 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6380 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6381 {
"objc_unsafeClaimAutoreleasedReturnValue",
6382 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6383 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6384 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6385 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6386 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6387 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6388 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6389 {
"objc_arc_annotation_topdown_bbstart",
6390 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6391 {
"objc_arc_annotation_topdown_bbend",
6392 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6393 {
"objc_arc_annotation_bottomup_bbstart",
6394 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6395 {
"objc_arc_annotation_bottomup_bbend",
6396 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6398 for (
auto &
I : RuntimeFuncs)
6399 UpgradeToIntrinsic(
I.first,
I.second);
6423 std::optional<bool> UseAddressDisc;
6426 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6427 for (
const MDNode *Flag : ModFlags->operands()) {
6429 if (Name && (*Name ==
"ptrauth-init-fini" ||
6430 *Name ==
"ptrauth-init-fini-address-discrimination"))
6435 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6436 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6437 constexpr unsigned ExpectedAddressMarker = 1;
6440 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6443 bool HasAddressDisc;
6444 if (!CPA->hasAddressDiscriminator())
6445 HasAddressDisc =
false;
6446 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6447 HasAddressDisc =
true;
6451 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6454 UseAddressDisc = HasAddressDisc;
6455 return CPA->getPointer();
6459 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6462 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6464 if (!GV || !GV->hasInitializer())
6468 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6471 std::vector<Constant *> NewStructors;
6472 NewStructors.reserve(OldStructorsArray->getNumOperands());
6474 for (
Use &U : OldStructorsArray->operands()) {
6483 Func = UpgradeSinglePointer(Func);
6487 NewStructors.push_back(
6496 if (GlobalArraysToUpgrade.
empty())
6498 assert(UseAddressDisc.has_value());
6500 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6501 GV->setInitializer(NewInit);
6504 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6514 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6518 bool HasObjCFlag =
false, HasClassProperties =
false;
6519 bool HasSwiftVersionFlag =
false;
6520 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6527 if (
Op->getNumOperands() != 3)
6541 if (ID->getString() ==
"Objective-C Image Info Version")
6543 if (ID->getString() ==
"Objective-C Class Properties")
6544 HasClassProperties =
true;
6546 if (ID->getString() ==
"PIC Level") {
6547 if (
auto *Behavior =
6549 uint64_t V = Behavior->getLimitedValue();
6555 if (ID->getString() ==
"PIE Level")
6556 if (
auto *Behavior =
6563 if (ID->getString() ==
"branch-target-enforcement" ||
6564 ID->getString().starts_with(
"sign-return-address")) {
6565 if (
auto *Behavior =
6571 Op->getOperand(1),
Op->getOperand(2)};
6581 if (ID->getString() ==
"Objective-C Image Info Section") {
6584 Value->getString().split(ValueComp,
" ");
6585 if (ValueComp.
size() != 1) {
6586 std::string NewValue;
6587 for (
auto &S : ValueComp)
6588 NewValue += S.str();
6599 if (ID->getString() ==
"Objective-C Garbage Collection") {
6602 assert(Md->getValue() &&
"Expected non-empty metadata");
6603 auto Type = Md->getValue()->getType();
6606 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
6607 if ((Val & 0xff) != Val) {
6608 HasSwiftVersionFlag =
true;
6609 SwiftABIVersion = (Val & 0xff00) >> 8;
6610 SwiftMajorVersion = (Val & 0xff000000) >> 24;
6611 SwiftMinorVersion = (Val & 0xff0000) >> 16;
6622 if (ID->getString() ==
"amdgpu_code_object_version") {
6625 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
6634 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
6663 if (HasObjCFlag && !HasClassProperties) {
6669 if (HasSwiftVersionFlag) {
6673 ConstantInt::get(Int8Ty, SwiftMajorVersion));
6675 ConstantInt::get(Int8Ty, SwiftMinorVersion));
6683 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
6687 auto MatchesVersion = [](
const MDNode *
Op) {
6688 return Op->getNumOperands() >= 3 &&
6702 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
6703 assert(
Op->getNumOperands() >= 2 &&
6704 "Expected at least 2 operands - name and linkage type");
6716 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
6727 auto TrimSpaces = [](
StringRef Section) -> std::string {
6729 Section.split(Components,
',');
6734 for (
auto Component : Components)
6735 OS <<
',' << Component.trim();
6740 for (
auto &GV : M.globals()) {
6741 if (!GV.hasSection())
6746 if (!Section.starts_with(
"__DATA, __objc_catlist"))
6751 GV.setSection(TrimSpaces(Section));
6767struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
6768 StrictFPUpgradeVisitor() =
default;
6771 if (!
Call.isStrictFP())
6777 Call.removeFnAttr(Attribute::StrictFP);
6778 Call.addFnAttr(Attribute::NoBuiltin);
6783struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
6784 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
6785 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
6787 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
6802 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
6803 StrictFPUpgradeVisitor SFPV;
6808 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
6809 F.getReturnType(),
F.getAttributes().getRetAttrs()));
6810 for (
auto &Arg :
F.args())
6812 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
6814 bool AddingAttrs =
false, RemovingAttrs =
false;
6815 AttrBuilder AttrsToAdd(
F.getContext());
6820 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
6821 A.isValid() &&
A.isStringAttribute()) {
6822 F.setSection(
A.getValueAsString());
6824 RemovingAttrs =
true;
6828 A.isValid() &&
A.isStringAttribute()) {
6831 AddingAttrs = RemovingAttrs =
true;
6834 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
6835 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
6837 RemovingAttrs =
true;
6838 if (
A.getValueAsString() ==
"true") {
6839 AttrsToAdd.addAttribute(
"uniform-work-group-size");
6848 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
6851 if (
A.getValueAsBool()) {
6852 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
6858 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
6859 RemovingAttrs =
true;
6866 bool HandleDenormalMode =
false;
6868 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
6871 DenormalFPMath = ParsedMode;
6873 AddingAttrs = RemovingAttrs =
true;
6874 HandleDenormalMode =
true;
6878 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
6882 DenormalFPMathF32 = ParsedMode;
6884 AddingAttrs = RemovingAttrs =
true;
6885 HandleDenormalMode =
true;
6889 if (HandleDenormalMode)
6890 AttrsToAdd.addDenormalFPEnvAttr(
6894 F.removeFnAttrs(AttrsToRemove);
6897 F.addFnAttrs(AttrsToAdd);
6903 if (!
F.hasFnAttribute(FnAttrName))
6904 F.addFnAttr(FnAttrName,
Value);
6911 if (!
F.hasFnAttribute(FnAttrName)) {
6913 F.addFnAttr(FnAttrName);
6915 auto A =
F.getFnAttribute(FnAttrName);
6916 if (
"false" ==
A.getValueAsString())
6917 F.removeFnAttr(FnAttrName);
6918 else if (
"true" ==
A.getValueAsString()) {
6919 F.removeFnAttr(FnAttrName);
6920 F.addFnAttr(FnAttrName);
6926 Triple T(M.getTargetTriple());
6927 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
6937 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6941 if (
Op->getNumOperands() != 3)
6950 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
6951 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
6952 : IDStr ==
"guarded-control-stack" ? &GCSValue
6953 : IDStr ==
"sign-return-address" ? &SRAValue
6954 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
6955 : IDStr ==
"sign-return-address-with-bkey"
6961 *ValPtr = CI->getZExtValue();
6967 bool BTE = BTEValue == 1;
6968 bool BPPLR = BPPLRValue == 1;
6969 bool GCS = GCSValue == 1;
6970 bool SRA = SRAValue == 1;
6973 if (SRA && SRAALLValue == 1)
6974 SignTypeValue =
"all";
6977 if (SRA && SRABKeyValue == 1)
6978 SignKeyValue =
"b_key";
6980 for (
Function &
F : M.getFunctionList()) {
6981 if (
F.isDeclaration())
6988 if (
auto A =
F.getFnAttribute(
"sign-return-address");
6989 A.isValid() &&
"none" ==
A.getValueAsString()) {
6990 F.removeFnAttr(
"sign-return-address");
6991 F.removeFnAttr(
"sign-return-address-key");
7007 if (SRAALLValue == 1)
7009 if (SRABKeyValue == 1)
7018 if (
T->getNumOperands() != 2)
7021 if (!Tag || Tag->getString() !=
"llvm.loop.distribute.enable")
7030 if (
T->getNumOperands() != 2)
7033 if (!Tag || Tag->getString() !=
"llvm.loop.vectorize.enable")
7044 :
"llvm.loop.vectorize.disable")});
7051 if (
T->getNumOperands() < 1)
7056 if (S->getString().starts_with(
"llvm.vectorizer."))
7062 StringRef OldPrefix =
"llvm.vectorizer.";
7065 if (OldTag ==
"llvm.vectorizer.unroll")
7077 if (
T->getNumOperands() < 1)
7090 :
"llvm.loop.distribute.disable")});
7098 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7106 if (NewTag->
getString() ==
"llvm.loop.vectorize.enable" &&
7111 Ops.reserve(
T->getNumOperands());
7112 Ops.push_back(NewTag);
7113 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7114 Ops.push_back(
T->getOperand(
I));
7131 if (
T->isDistinct()) {
7132 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7144 Ops.reserve(
T->getNumOperands());
7155 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7156 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7157 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7160 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7162 auto I =
DL.find(
"-n64-");
7164 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7169 std::string Res =
DL.str();
7172 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7173 Res.append(Res.empty() ?
"G1" :
"-G1");
7181 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7182 Res.append(
"-ni:7:8:9");
7184 if (
DL.ends_with(
"ni:7"))
7186 if (
DL.ends_with(
"ni:7:8"))
7191 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7192 Res.append(
"-p7:160:256:256:32");
7193 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7194 Res.append(
"-p8:128:128:128:48");
7195 constexpr StringRef OldP8(
"-p8:128:128-");
7196 if (
DL.contains(OldP8))
7197 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7198 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7199 Res.append(
"-p9:192:256:256:32");
7203 if (!
DL.contains(
"m:e"))
7204 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7209 if (
T.isSystemZ() && !
DL.empty()) {
7211 if (!
DL.contains(
"-S64"))
7212 return "E-S64" +
DL.drop_front(1).str();
7216 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7219 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7220 if (!
DL.contains(AddrSpaces)) {
7222 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7223 if (R.match(Res, &
Groups))
7229 if (
T.isAArch64()) {
7231 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7232 Res.append(
"-Fn32");
7233 AddPtr32Ptr64AddrSpaces();
7237 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7241 std::string I64 =
"-i64:64";
7242 std::string I128 =
"-i128:128";
7244 size_t Pos = Res.find(I64);
7245 if (Pos !=
size_t(-1))
7246 Res.insert(Pos + I64.size(), I128);
7250 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7251 size_t Pos = Res.find(
"-S128");
7254 Res.insert(Pos,
"-f64:32:64");
7260 AddPtr32Ptr64AddrSpaces();
7268 if (!
T.isOSIAMCU()) {
7269 std::string I128 =
"-i128:128";
7272 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7273 if (R.match(Res, &
Groups))
7281 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7283 auto I =
Ref.find(
"-f80:32-");
7285 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7293 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7296 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7297 B.removeAttribute(
"no-frame-pointer-elim");
7299 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7301 if (FramePointer !=
"all")
7302 FramePointer =
"non-leaf";
7303 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7305 if (!FramePointer.
empty())
7306 B.addAttribute(
"frame-pointer", FramePointer);
7308 A =
B.getAttribute(
"null-pointer-is-valid");
7311 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7312 B.removeAttribute(
"null-pointer-is-valid");
7313 if (NullPointerIsValid)
7314 B.addAttribute(Attribute::NullPointerIsValid);
7317 A =
B.getAttribute(
"uniform-work-group-size");
7321 bool IsTrue = Val ==
"true";
7322 B.removeAttribute(
"uniform-work-group-size");
7324 B.addAttribute(
"uniform-work-group-size");
7335 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 bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static Metadata * makeVectorizeEnableNode(LLVMContext &C, const MDOperand &Op)
Build the single-operand vectorize enable/disable node that replaces a boolean operand: nonzero -> en...
static Metadata * upgradeLoopArgument(Metadata *MD)
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 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 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 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 Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool isOldDistributeEnable(const MDTuple *T)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
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 bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
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 isOldVectorizeEnable(const MDTuple *T)
Old two-operand form: !
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)
static cl::opt< bool > DisableAutoUpgradeDebugInfo("disable-auto-upgrade-debug-info", cl::desc("Disable autoupgrade of debug info"))
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 Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
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 void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
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 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.
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.
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ ICMP_ULT
unsigned less than
@ ICMP_SGE
signed greater or equal
@ ICMP_ULE
unsigned less or equal
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.
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',...
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
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
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...
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.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
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 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 ID lookupIntrinsicID(StringRef Name)
This does the actual lookup of an intrinsic ID which matches the given function name.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
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)
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_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.
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.
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
@ Dynamic
Denotes mode unknown at compile time.
ArrayRef(const T &OneElt) -> ArrayRef< T >
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...
@ Default
The result value is uniform if and only if all operands are uniform.
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.