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"
66 cl::desc(
"Disable autoupgrade of debug info"));
85 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
100 Type *LastArgType =
F->getFunctionType()->getParamType(
101 F->getFunctionType()->getNumParams() - 1);
116 if (
F->getReturnType()->isVectorTy())
129 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
130 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
147 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
148 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
162 if (
F->getReturnType()->getScalarType()->isBFloatTy())
172 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
186 if (Name.consume_front(
"avx."))
187 return (Name.starts_with(
"blend.p") ||
188 Name ==
"cvt.ps2.pd.256" ||
189 Name ==
"cvtdq2.pd.256" ||
190 Name ==
"cvtdq2.ps.256" ||
191 Name.starts_with(
"movnt.") ||
192 Name.starts_with(
"sqrt.p") ||
193 Name.starts_with(
"storeu.") ||
194 Name.starts_with(
"vbroadcast.s") ||
195 Name.starts_with(
"vbroadcastf128") ||
196 Name.starts_with(
"vextractf128.") ||
197 Name.starts_with(
"vinsertf128.") ||
198 Name.starts_with(
"vperm2f128.") ||
199 Name.starts_with(
"vpermil."));
201 if (Name.consume_front(
"avx2."))
202 return (Name ==
"movntdqa" ||
203 Name.starts_with(
"pabs.") ||
204 Name.starts_with(
"padds.") ||
205 Name.starts_with(
"paddus.") ||
206 Name.starts_with(
"pblendd.") ||
208 Name.starts_with(
"pbroadcast") ||
209 Name.starts_with(
"pcmpeq.") ||
210 Name.starts_with(
"pcmpgt.") ||
211 Name.starts_with(
"pmax") ||
212 Name.starts_with(
"pmin") ||
213 Name.starts_with(
"pmovsx") ||
214 Name.starts_with(
"pmovzx") ||
215 Name.starts_with(
"pmulh.w") ||
216 Name.starts_with(
"pmulhu.w") ||
218 Name ==
"pmulu.dq" ||
219 Name.starts_with(
"psll.dq") ||
220 Name.starts_with(
"psrl.dq") ||
221 Name.starts_with(
"psubs.") ||
222 Name.starts_with(
"psubus.") ||
223 Name.starts_with(
"vbroadcast") ||
224 Name ==
"vbroadcasti128" ||
225 Name ==
"vextracti128" ||
226 Name ==
"vinserti128" ||
227 Name ==
"vperm2i128");
229 if (Name.consume_front(
"avx512.")) {
230 if (Name.consume_front(
"mask."))
232 return (Name.starts_with(
"add.p") ||
233 Name.starts_with(
"and.") ||
234 Name.starts_with(
"andn.") ||
235 Name.starts_with(
"broadcast.s") ||
236 Name.starts_with(
"broadcastf32x4.") ||
237 Name.starts_with(
"broadcastf32x8.") ||
238 Name.starts_with(
"broadcastf64x2.") ||
239 Name.starts_with(
"broadcastf64x4.") ||
240 Name.starts_with(
"broadcasti32x4.") ||
241 Name.starts_with(
"broadcasti32x8.") ||
242 Name.starts_with(
"broadcasti64x2.") ||
243 Name.starts_with(
"broadcasti64x4.") ||
244 Name.starts_with(
"cmp.b") ||
245 Name.starts_with(
"cmp.d") ||
246 Name.starts_with(
"cmp.q") ||
247 Name.starts_with(
"cmp.w") ||
248 Name.starts_with(
"compress.b") ||
249 Name.starts_with(
"compress.d") ||
250 Name.starts_with(
"compress.p") ||
251 Name.starts_with(
"compress.q") ||
252 Name.starts_with(
"compress.store.") ||
253 Name.starts_with(
"compress.w") ||
254 Name.starts_with(
"conflict.") ||
255 Name.starts_with(
"cvtdq2pd.") ||
256 Name.starts_with(
"cvtdq2ps.") ||
257 Name ==
"cvtpd2dq.256" ||
258 Name ==
"cvtpd2ps.256" ||
259 Name ==
"cvtps2pd.128" ||
260 Name ==
"cvtps2pd.256" ||
261 Name.starts_with(
"cvtqq2pd.") ||
262 Name ==
"cvtqq2ps.256" ||
263 Name ==
"cvtqq2ps.512" ||
264 Name ==
"cvttpd2dq.256" ||
265 Name ==
"cvttps2dq.128" ||
266 Name ==
"cvttps2dq.256" ||
267 Name.starts_with(
"cvtudq2pd.") ||
268 Name.starts_with(
"cvtudq2ps.") ||
269 Name.starts_with(
"cvtuqq2pd.") ||
270 Name ==
"cvtuqq2ps.256" ||
271 Name ==
"cvtuqq2ps.512" ||
272 Name.starts_with(
"dbpsadbw.") ||
273 Name.starts_with(
"div.p") ||
274 Name.starts_with(
"expand.b") ||
275 Name.starts_with(
"expand.d") ||
276 Name.starts_with(
"expand.load.") ||
277 Name.starts_with(
"expand.p") ||
278 Name.starts_with(
"expand.q") ||
279 Name.starts_with(
"expand.w") ||
280 Name.starts_with(
"fpclass.p") ||
281 Name.starts_with(
"insert") ||
282 Name.starts_with(
"load.") ||
283 Name.starts_with(
"loadu.") ||
284 Name.starts_with(
"lzcnt.") ||
285 Name.starts_with(
"max.p") ||
286 Name.starts_with(
"min.p") ||
287 Name.starts_with(
"movddup") ||
288 Name.starts_with(
"move.s") ||
289 Name.starts_with(
"movshdup") ||
290 Name.starts_with(
"movsldup") ||
291 Name.starts_with(
"mul.p") ||
292 Name.starts_with(
"or.") ||
293 Name.starts_with(
"pabs.") ||
294 Name.starts_with(
"packssdw.") ||
295 Name.starts_with(
"packsswb.") ||
296 Name.starts_with(
"packusdw.") ||
297 Name.starts_with(
"packuswb.") ||
298 Name.starts_with(
"padd.") ||
299 Name.starts_with(
"padds.") ||
300 Name.starts_with(
"paddus.") ||
301 Name.starts_with(
"palignr.") ||
302 Name.starts_with(
"pand.") ||
303 Name.starts_with(
"pandn.") ||
304 Name.starts_with(
"pavg") ||
305 Name.starts_with(
"pbroadcast") ||
306 Name.starts_with(
"pcmpeq.") ||
307 Name.starts_with(
"pcmpgt.") ||
308 Name.starts_with(
"perm.df.") ||
309 Name.starts_with(
"perm.di.") ||
310 Name.starts_with(
"permvar.") ||
311 Name.starts_with(
"pmaddubs.w.") ||
312 Name.starts_with(
"pmaddw.d.") ||
313 Name.starts_with(
"pmax") ||
314 Name.starts_with(
"pmin") ||
315 Name ==
"pmov.qd.256" ||
316 Name ==
"pmov.qd.512" ||
317 Name ==
"pmov.wb.256" ||
318 Name ==
"pmov.wb.512" ||
319 Name.starts_with(
"pmovsx") ||
320 Name.starts_with(
"pmovzx") ||
321 Name.starts_with(
"pmul.dq.") ||
322 Name.starts_with(
"pmul.hr.sw.") ||
323 Name.starts_with(
"pmulh.w.") ||
324 Name.starts_with(
"pmulhu.w.") ||
325 Name.starts_with(
"pmull.") ||
326 Name.starts_with(
"pmultishift.qb.") ||
327 Name.starts_with(
"pmulu.dq.") ||
328 Name.starts_with(
"por.") ||
329 Name.starts_with(
"prol.") ||
330 Name.starts_with(
"prolv.") ||
331 Name.starts_with(
"pror.") ||
332 Name.starts_with(
"prorv.") ||
333 Name.starts_with(
"pshuf.b.") ||
334 Name.starts_with(
"pshuf.d.") ||
335 Name.starts_with(
"pshufh.w.") ||
336 Name.starts_with(
"pshufl.w.") ||
337 Name.starts_with(
"psll.d") ||
338 Name.starts_with(
"psll.q") ||
339 Name.starts_with(
"psll.w") ||
340 Name.starts_with(
"pslli") ||
341 Name.starts_with(
"psllv") ||
342 Name.starts_with(
"psra.d") ||
343 Name.starts_with(
"psra.q") ||
344 Name.starts_with(
"psra.w") ||
345 Name.starts_with(
"psrai") ||
346 Name.starts_with(
"psrav") ||
347 Name.starts_with(
"psrl.d") ||
348 Name.starts_with(
"psrl.q") ||
349 Name.starts_with(
"psrl.w") ||
350 Name.starts_with(
"psrli") ||
351 Name.starts_with(
"psrlv") ||
352 Name.starts_with(
"psub.") ||
353 Name.starts_with(
"psubs.") ||
354 Name.starts_with(
"psubus.") ||
355 Name.starts_with(
"pternlog.") ||
356 Name.starts_with(
"punpckh") ||
357 Name.starts_with(
"punpckl") ||
358 Name.starts_with(
"pxor.") ||
359 Name.starts_with(
"shuf.f") ||
360 Name.starts_with(
"shuf.i") ||
361 Name.starts_with(
"shuf.p") ||
362 Name.starts_with(
"sqrt.p") ||
363 Name.starts_with(
"store.b.") ||
364 Name.starts_with(
"store.d.") ||
365 Name.starts_with(
"store.p") ||
366 Name.starts_with(
"store.q.") ||
367 Name.starts_with(
"store.w.") ||
368 Name ==
"store.ss" ||
369 Name.starts_with(
"storeu.") ||
370 Name.starts_with(
"sub.p") ||
371 Name.starts_with(
"ucmp.") ||
372 Name.starts_with(
"unpckh.") ||
373 Name.starts_with(
"unpckl.") ||
374 Name.starts_with(
"valign.") ||
375 Name ==
"vcvtph2ps.128" ||
376 Name ==
"vcvtph2ps.256" ||
377 Name.starts_with(
"vextract") ||
378 Name.starts_with(
"vfmadd.") ||
379 Name.starts_with(
"vfmaddsub.") ||
380 Name.starts_with(
"vfnmadd.") ||
381 Name.starts_with(
"vfnmsub.") ||
382 Name.starts_with(
"vpdpbusd.") ||
383 Name.starts_with(
"vpdpbusds.") ||
384 Name.starts_with(
"vpdpwssd.") ||
385 Name.starts_with(
"vpdpwssds.") ||
386 Name.starts_with(
"vpermi2var.") ||
387 Name.starts_with(
"vpermil.p") ||
388 Name.starts_with(
"vpermilvar.") ||
389 Name.starts_with(
"vpermt2var.") ||
390 Name.starts_with(
"vpmadd52") ||
391 Name.starts_with(
"vpshld.") ||
392 Name.starts_with(
"vpshldv.") ||
393 Name.starts_with(
"vpshrd.") ||
394 Name.starts_with(
"vpshrdv.") ||
395 Name.starts_with(
"vpshufbitqmb.") ||
396 Name.starts_with(
"xor."));
398 if (Name.consume_front(
"mask3."))
400 return (Name.starts_with(
"vfmadd.") ||
401 Name.starts_with(
"vfmaddsub.") ||
402 Name.starts_with(
"vfmsub.") ||
403 Name.starts_with(
"vfmsubadd.") ||
404 Name.starts_with(
"vfnmsub."));
406 if (Name.consume_front(
"maskz."))
408 return (Name.starts_with(
"pternlog.") ||
409 Name.starts_with(
"vfmadd.") ||
410 Name.starts_with(
"vfmaddsub.") ||
411 Name.starts_with(
"vpdpbusd.") ||
412 Name.starts_with(
"vpdpbusds.") ||
413 Name.starts_with(
"vpdpwssd.") ||
414 Name.starts_with(
"vpdpwssds.") ||
415 Name.starts_with(
"vpermt2var.") ||
416 Name.starts_with(
"vpmadd52") ||
417 Name.starts_with(
"vpshldv.") ||
418 Name.starts_with(
"vpshrdv."));
421 return (Name ==
"movntdqa" ||
422 Name ==
"pmul.dq.512" ||
423 Name ==
"pmulu.dq.512" ||
424 Name.starts_with(
"broadcastm") ||
425 Name.starts_with(
"cmp.p") ||
426 Name.starts_with(
"cvtb2mask.") ||
427 Name.starts_with(
"cvtd2mask.") ||
428 Name.starts_with(
"cvtmask2") ||
429 Name.starts_with(
"cvtq2mask.") ||
430 Name ==
"cvtusi2sd" ||
431 Name.starts_with(
"cvtw2mask.") ||
436 Name ==
"kortestc.w" ||
437 Name ==
"kortestz.w" ||
438 Name.starts_with(
"kunpck") ||
441 Name.starts_with(
"padds.") ||
442 Name.starts_with(
"pbroadcast") ||
443 Name.starts_with(
"pmulh.w") ||
444 Name.starts_with(
"pmulhu.w") ||
445 Name.starts_with(
"prol") ||
446 Name.starts_with(
"pror") ||
447 Name.starts_with(
"psll.dq") ||
448 Name.starts_with(
"psrl.dq") ||
449 Name.starts_with(
"psubs.") ||
450 Name.starts_with(
"ptestm") ||
451 Name.starts_with(
"ptestnm") ||
452 Name.starts_with(
"storent.") ||
453 Name.starts_with(
"vbroadcast.s") ||
454 Name.starts_with(
"vpshld.") ||
455 Name.starts_with(
"vpshrd."));
458 if (Name.consume_front(
"fma."))
459 return (Name.starts_with(
"vfmadd.") ||
460 Name.starts_with(
"vfmsub.") ||
461 Name.starts_with(
"vfmsubadd.") ||
462 Name.starts_with(
"vfnmadd.") ||
463 Name.starts_with(
"vfnmsub."));
465 if (Name.consume_front(
"fma4."))
466 return Name.starts_with(
"vfmadd.s");
468 if (Name.consume_front(
"sse."))
469 return (Name ==
"add.ss" ||
470 Name ==
"cvtsi2ss" ||
471 Name ==
"cvtsi642ss" ||
474 Name.starts_with(
"sqrt.p") ||
476 Name.starts_with(
"storeu.") ||
479 if (Name.consume_front(
"sse2."))
480 return (Name ==
"add.sd" ||
481 Name ==
"cvtdq2pd" ||
482 Name ==
"cvtdq2ps" ||
483 Name ==
"cvtps2pd" ||
484 Name ==
"cvtsi2sd" ||
485 Name ==
"cvtsi642sd" ||
486 Name ==
"cvtss2sd" ||
489 Name.starts_with(
"padds.") ||
490 Name.starts_with(
"paddus.") ||
491 Name.starts_with(
"pcmpeq.") ||
492 Name.starts_with(
"pcmpgt.") ||
498 Name ==
"pmulhu.w" ||
499 Name ==
"pmulu.dq" ||
500 Name.starts_with(
"pshuf") ||
501 Name.starts_with(
"psll.dq") ||
502 Name.starts_with(
"psrl.dq") ||
503 Name.starts_with(
"psubs.") ||
504 Name.starts_with(
"psubus.") ||
505 Name.starts_with(
"sqrt.p") ||
507 Name ==
"storel.dq" ||
508 Name.starts_with(
"storeu.") ||
511 if (Name.consume_front(
"sse41."))
512 return (Name.starts_with(
"blendp") ||
513 Name ==
"movntdqa" ||
523 Name.starts_with(
"pmovsx") ||
524 Name.starts_with(
"pmovzx") ||
527 if (Name.consume_front(
"sse42."))
528 return Name ==
"crc32.64.8";
530 if (Name.consume_front(
"sse4a."))
531 return Name.starts_with(
"movnt.");
533 if (Name.consume_front(
"ssse3."))
534 return (Name ==
"pabs.b.128" ||
535 Name ==
"pabs.d.128" ||
536 Name ==
"pabs.w.128");
538 if (Name.consume_front(
"xop."))
539 return (Name ==
"vpcmov" ||
540 Name ==
"vpcmov.256" ||
541 Name.starts_with(
"vpcom") ||
542 Name.starts_with(
"vprot"));
544 if (Name.consume_front(
"bmi."))
545 return (Name.starts_with(
"pdep.") ||
546 Name.starts_with(
"pext."));
548 return (Name ==
"addcarry.u32" ||
549 Name ==
"addcarry.u64" ||
550 Name ==
"addcarryx.u32" ||
551 Name ==
"addcarryx.u64" ||
552 Name ==
"subborrow.u32" ||
553 Name ==
"subborrow.u64" ||
554 Name.starts_with(
"vcvtph2ps."));
560 if (!Name.consume_front(
"x86."))
568 if (Name ==
"rdtscp") {
570 if (
F->getFunctionType()->getNumParams() == 0)
575 Intrinsic::x86_rdtscp);
582 if (Name.consume_front(
"sse41.ptest")) {
584 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
585 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
586 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
599 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
600 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
601 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
602 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
603 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
604 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
609 if (Name.consume_front(
"avx512.")) {
610 if (Name.consume_front(
"mask.cmp.")) {
613 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
614 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
615 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
616 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
617 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
618 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
622 }
else if (Name.starts_with(
"vpdpbusd.") ||
623 Name.starts_with(
"vpdpbusds.")) {
626 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
627 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
628 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
629 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
630 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
631 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
635 }
else if (Name.starts_with(
"vpdpwssd.") ||
636 Name.starts_with(
"vpdpwssds.")) {
639 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
640 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
641 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
642 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
643 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
644 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
652 if (Name.consume_front(
"avx2.")) {
653 if (Name.consume_front(
"vpdpb")) {
656 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
657 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
658 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
659 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
660 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
661 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
662 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
663 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
664 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
665 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
666 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
667 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
671 }
else if (Name.consume_front(
"vpdpw")) {
674 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
675 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
676 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
677 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
678 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
679 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
680 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
681 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
682 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
683 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
684 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
685 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
693 if (Name.consume_front(
"avx10.")) {
694 if (Name.consume_front(
"vpdpb")) {
697 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
698 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
699 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
700 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
701 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
702 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
706 }
else if (Name.consume_front(
"vpdpw")) {
708 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
709 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
710 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
711 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
712 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
713 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
721 if (Name.consume_front(
"avx512bf16.")) {
724 .
Case(
"cvtne2ps2bf16.128",
725 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
726 .
Case(
"cvtne2ps2bf16.256",
727 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
728 .
Case(
"cvtne2ps2bf16.512",
729 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
730 .
Case(
"mask.cvtneps2bf16.128",
731 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
732 .
Case(
"cvtneps2bf16.256",
733 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
734 .
Case(
"cvtneps2bf16.512",
735 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
742 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
743 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
744 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
751 if (Name.consume_front(
"xop.")) {
753 if (Name.starts_with(
"vpermil2")) {
756 auto Idx =
F->getFunctionType()->getParamType(2);
757 if (Idx->isFPOrFPVectorTy()) {
758 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
759 unsigned EltSize = Idx->getScalarSizeInBits();
760 if (EltSize == 64 && IdxSize == 128)
761 ID = Intrinsic::x86_xop_vpermil2pd;
762 else if (EltSize == 32 && IdxSize == 128)
763 ID = Intrinsic::x86_xop_vpermil2ps;
764 else if (EltSize == 64 && IdxSize == 256)
765 ID = Intrinsic::x86_xop_vpermil2pd_256;
767 ID = Intrinsic::x86_xop_vpermil2ps_256;
769 }
else if (
F->arg_size() == 2)
772 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
773 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
784 if (Name ==
"seh.recoverfp") {
786 Intrinsic::eh_recoverfp);
798 if (Name.starts_with(
"rbit")) {
801 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
805 if (Name ==
"thread.pointer") {
808 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
812 bool Neon = Name.consume_front(
"neon.");
817 if (Name.consume_front(
"bfdot.")) {
821 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
826 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
827 assert((OperandWidth == 64 || OperandWidth == 128) &&
828 "Unexpected operand width");
830 std::array<Type *, 2> Tys{
841 if (Name.consume_front(
"bfm")) {
843 if (Name.consume_back(
".v4f32.v16i8")) {
889 F->arg_begin()->getType());
893 if (Name.consume_front(
"vst")) {
895 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
899 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
900 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
903 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
904 Intrinsic::arm_neon_vst4lane};
906 auto fArgs =
F->getFunctionType()->params();
907 Type *Tys[] = {fArgs[0], fArgs[1]};
910 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
913 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
922 if (Name.consume_front(
"mve.")) {
924 if (Name ==
"vctp64") {
934 if (Name.starts_with(
"vrintn.v")) {
936 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
941 if (Name.consume_back(
".v4i1")) {
943 if (Name.consume_back(
".predicated.v2i64.v4i32"))
945 return Name ==
"mull.int" || Name ==
"vqdmull";
947 if (Name.consume_back(
".v2i64")) {
949 bool IsGather = Name.consume_front(
"vldr.gather.");
950 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
951 if (Name.consume_front(
"base.")) {
953 Name.consume_front(
"wb.");
956 return Name ==
"predicated.v2i64";
959 if (Name.consume_front(
"offset.predicated."))
960 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
961 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
974 if (Name.consume_front(
"cde.vcx")) {
976 if (Name.consume_back(
".predicated.v2i64.v4i1"))
978 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
979 Name ==
"3q" || Name ==
"3qa";
993 F->arg_begin()->getType());
999 .
Case(
"smax", Intrinsic::smax)
1000 .
Case(
"smin", Intrinsic::smin)
1001 .
Case(
"umax", Intrinsic::umax)
1002 .
Case(
"umin", Intrinsic::umin)
1005 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1008 F->getReturnType());
1012 if (Name.starts_with(
"addp")) {
1014 if (
F->arg_size() != 2)
1017 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1019 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1025 if (Name.starts_with(
"bfcvt")) {
1031 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1038 if (Name.consume_front(
"sve.")) {
1040 if (Name.consume_front(
"bf")) {
1041 if (Name ==
"mmla") {
1042 Type *Tys[] = {
F->getReturnType(),
1043 std::next(
F->arg_begin())->getType()};
1045 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1048 if (Name.consume_back(
".lane")) {
1052 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1053 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1054 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1066 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1071 if (Name.consume_front(
"convert.from.svbool")) {
1074 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1077 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1082 if (Name.consume_front(
"convert.to.svbool")) {
1085 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1088 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1093 if (Name.consume_front(
"addqv")) {
1095 if (!
F->getReturnType()->isFPOrFPVectorTy())
1098 auto Args =
F->getFunctionType()->params();
1099 Type *Tys[] = {
F->getReturnType(), Args[1]};
1101 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1105 if (Name.consume_front(
"ld")) {
1107 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1108 if (LdRegex.
match(Name)) {
1114 "Expected 2 arguments for ld* intrinsic.");
1115 Type *PtrTy =
F->getArg(1)->getType();
1118 Intrinsic::aarch64_sve_ld2_sret,
1119 Intrinsic::aarch64_sve_ld3_sret,
1120 Intrinsic::aarch64_sve_ld4_sret,
1123 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1129 if (Name.consume_front(
"tuple.")) {
1131 if (Name.starts_with(
"get")) {
1133 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1135 F->getParent(), Intrinsic::vector_extract, Tys);
1139 if (Name.starts_with(
"set")) {
1141 auto Args =
F->getFunctionType()->params();
1142 Type *Tys[] = {Args[0], Args[2], Args[1]};
1144 F->getParent(), Intrinsic::vector_insert, Tys);
1148 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1149 if (CreateTupleRegex.
match(Name)) {
1151 auto Args =
F->getFunctionType()->params();
1152 Type *Tys[] = {
F->getReturnType(), Args[1]};
1154 F->getParent(), Intrinsic::vector_insert, Tys);
1160 if (Name.starts_with(
"rev.nxv")) {
1163 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1169 if (Name.consume_front(
"sme.")) {
1171 if (Name.consume_front(
"ftmopa.")) {
1176 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1177 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1195#define NVVM_TMA_G2S_MODES(M) \
1196 M(tile_1d, "tile.1d") \
1197 M(tile_2d, "tile.2d") \
1198 M(tile_3d, "tile.3d") \
1199 M(tile_4d, "tile.4d") \
1200 M(tile_5d, "tile.5d") \
1201 M(tile_gather4_2d, "tile.gather4.2d") \
1202 M(im2col_3d, "im2col.3d") \
1203 M(im2col_4d, "im2col.4d") \
1204 M(im2col_5d, "im2col.5d") \
1205 M(im2col_w_3d, "im2col.w.3d") \
1206 M(im2col_w_4d, "im2col.w.4d") \
1207 M(im2col_w_5d, "im2col.w.5d") \
1208 M(im2col_w_128_3d, "im2col.w.128.3d") \
1209 M(im2col_w_128_4d, "im2col.w.128.4d") \
1210 M(im2col_w_128_5d, "im2col.w.128.5d")
1222 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1225#define G2S_ID(ID_SUFFIX, NAME) \
1226 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1236 size_t NumParams =
F->getFunctionType()->getNumParams();
1240 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1247 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1248 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1249 "expected the i64 cache-hint after the multicast mask");
1250 Type *MaskTy = Params[MaskIdx];
1265 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1268#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1269 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1281 if (!
F->getFunctionType()
1282 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1298 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1303 size_t NumParams =
F->getFunctionType()->getNumParams();
1304 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1308 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1313 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1326 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1331 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1334 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1354 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1357 auto [RedOpName, ShapeName] = Name.split(
'.');
1362 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1363 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1364 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1365 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1366 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1367 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1368 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1369 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1375 if (Name.consume_front(
"mapa.shared.cluster"))
1376 if (
F->getReturnType()->getPointerAddressSpace() ==
1378 return Intrinsic::nvvm_mapa_shared_cluster;
1380 if (Name.consume_front(
"cp.async.bulk.")) {
1383 .
Case(
"shared.cta.to.cluster",
1384 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1388 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1398 if (!Name.consume_front(
"tcgen05.commit."))
1401 if (Name.consume_front(
"shared."))
1403 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1404 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1407 if (Name.consume_front(
"mc.shared.")) {
1409 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1413 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1414 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1423 if (
F->arg_size() != 2)
1426 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1427 Name.consume_front(
"tcgen05.alloc."))
1429 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1430 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1433 if (Name.consume_front(
"tcgen05.dealloc."))
1435 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1436 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1443 if (Name.consume_front(
"fma.rn."))
1445 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1446 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1447 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1448 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1451 if (Name.consume_front(
"fmax."))
1453 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1454 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1455 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1456 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1457 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1458 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1459 .
Case(
"ftz.nan.xorsign.abs.bf16",
1460 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1461 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1462 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1463 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1464 .
Case(
"ftz.xorsign.abs.bf16x2",
1465 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1466 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1467 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1468 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1469 .
Case(
"nan.xorsign.abs.bf16x2",
1470 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1471 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1472 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1475 if (Name.consume_front(
"fmin."))
1477 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1478 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1479 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1480 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1481 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1482 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1483 .
Case(
"ftz.nan.xorsign.abs.bf16",
1484 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1485 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1486 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1487 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1488 .
Case(
"ftz.xorsign.abs.bf16x2",
1489 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1490 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1491 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1492 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1493 .
Case(
"nan.xorsign.abs.bf16x2",
1494 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1495 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1496 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1499 if (Name.consume_front(
"neg."))
1501 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1502 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1511 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1516 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1519 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1522 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1523 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1529static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1531 auto [Modifiers,
Type] = Name.rsplit(
'.');
1533 return std::nullopt;
1543 return std::nullopt;
1546 .
Case(
"", Intrinsic::nvvm_fadd)
1547 .
Case(
".ftz", Intrinsic::nvvm_fadd_ftz)
1548 .
Case(
".sat", Intrinsic::nvvm_fadd_sat)
1549 .
Case(
".ftz.sat", Intrinsic::nvvm_fadd_ftz_sat)
1552 return std::nullopt;
1558 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1561 return Intrinsic::nvvm_mbarrier_init;
1565 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1566 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1567 Name.consume_front(
"param");
1571 if (!Name.consume_front(
"vp."))
1600 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1601 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1608 if (!Name.consume_front(
"vp."))
1628 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1629 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1634 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1646 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1657 if (Name.starts_with(
"to.fp16")) {
1661 FuncTy->getReturnType());
1664 if (Name.starts_with(
"from.fp16")) {
1668 FuncTy->getReturnType());
1678 if (Defaults.empty())
1681 unsigned FullArgCount = FirstDefault + Defaults.size();
1684 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1687 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1689 NumMissingTrailingParams))
1692 return FullArgCount;
1699 unsigned FullArgCount =
1701 if (FullArgCount == 0)
1707 "total number of default args does not match intrinsic signature");
1712 bool CanUpgradeDebugIntrinsicsToRecords) {
1713 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1718 if (!Name.consume_front(
"llvm.") || Name.empty())
1724 bool IsArm = Name.consume_front(
"arm.");
1725 if (IsArm || Name.consume_front(
"aarch64.")) {
1731 if (Name.consume_front(
"amdgcn.")) {
1732 if (Name ==
"alignbit") {
1735 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1739 if (Name.consume_front(
"atomic.")) {
1740 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1741 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1750 if (Name.starts_with(
"addrspacecast.nonnull")) {
1757 switch (
F->getIntrinsicID()) {
1761 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1762 if (
F->arg_size() == 7) {
1767 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1768 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1769 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1770 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1771 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1772 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1773 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1774 if (
F->arg_size() == 8) {
1781 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1782 Name.consume_front(
"flat.atomic.")) {
1783 if (Name.starts_with(
"fadd") ||
1785 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1786 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1794 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1799 if (Name.starts_with(
"ldexp.")) {
1802 F->getParent(), Intrinsic::ldexp,
1803 {F->getReturnType(), F->getArg(1)->getType()});
1812 if (
F->arg_size() == 1) {
1813 if (Name.consume_front(
"convert.")) {
1827 F->arg_begin()->getType());
1833 if (Name ==
"coro.end" &&
1834 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1835 CoroEndID = Intrinsic::coro_end;
1836 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1837 CoroEndID = Intrinsic::coro_end_async;
1848 if (Name.consume_front(
"dbg.")) {
1850 if (CanUpgradeDebugIntrinsicsToRecords) {
1851 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1852 Name ==
"declare" || Name ==
"label") {
1861 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1864 Intrinsic::dbg_value);
1871 if (Name.consume_front(
"experimental.vector.")) {
1877 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1878 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1879 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1880 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1881 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1883 Intrinsic::vector_partial_reduce_add)
1886 const auto *FT =
F->getFunctionType();
1888 if (ID == Intrinsic::vector_extract ||
1889 ID == Intrinsic::vector_interleave2)
1892 if (ID != Intrinsic::vector_interleave2)
1894 if (ID == Intrinsic::vector_insert ||
1895 ID == Intrinsic::vector_partial_reduce_add)
1903 if (Name.consume_front(
"reduce.")) {
1905 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1906 if (R.match(Name, &
Groups))
1908 .
Case(
"add", Intrinsic::vector_reduce_add)
1909 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1910 .
Case(
"and", Intrinsic::vector_reduce_and)
1911 .
Case(
"or", Intrinsic::vector_reduce_or)
1912 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1913 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1914 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1915 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1916 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1917 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1918 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1923 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1928 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1929 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1934 auto Args =
F->getFunctionType()->params();
1936 {Args[V2 ? 1 : 0]});
1942 if (Name.consume_front(
"splice"))
1946 if (Name.consume_front(
"experimental.stepvector.")) {
1950 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1955 if (Name.starts_with(
"flt.rounds")) {
1958 Intrinsic::get_rounding);
1963 if (Name.starts_with(
"invariant.group.barrier")) {
1965 auto Args =
F->getFunctionType()->params();
1966 Type* ObjectPtr[1] = {Args[0]};
1969 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1974 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1975 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1976 if (IsLifetimeStart || IsLifetimeEnd) {
1977 if (
F->arg_size() == 2) {
1978 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1979 : Intrinsic::lifetime_end;
1984 F->getArg(1)->getType());
1986 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2006 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2007 .StartsWith(
"memmove.", Intrinsic::memmove)
2009 if (
F->arg_size() == 5) {
2013 F->getFunctionType()->params().slice(0, 3);
2019 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2022 const auto *FT =
F->getFunctionType();
2023 Type *ParamTypes[2] = {
2024 FT->getParamType(0),
2028 Intrinsic::memset, ParamTypes);
2034 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2035 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2036 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2037 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2039 if (MaskedID &&
F->arg_size() == 4) {
2041 if (MaskedID == Intrinsic::masked_load ||
2042 MaskedID == Intrinsic::masked_gather) {
2044 F->getParent(), MaskedID,
2045 {F->getReturnType(), F->getArg(0)->getType()});
2049 F->getParent(), MaskedID,
2050 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2056 if (Name.consume_front(
"nvvm.")) {
2058 if (
F->arg_size() == 1) {
2061 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2062 .Case(
"clz.i", Intrinsic::ctlz)
2063 .
Case(
"popc.i", Intrinsic::ctpop)
2067 {F->getReturnType()});
2070 }
else if (
F->arg_size() == 2) {
2073 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2074 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2075 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2076 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2077 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2078 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2082 {F->getReturnType()});
2119 F->getParent(), IID,
F->getReturnType(),
2120 F->getFunctionType()->params());
2131 {F->getArg(0)->getType()});
2179 F->getArg(0)->getType());
2187 bool Expand =
false;
2188 if (Name.consume_front(
"abs."))
2191 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2192 else if (Name.consume_front(
"fabs."))
2194 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2195 else if (Name.consume_front(
"add."))
2198 else if (Name.consume_front(
"ex2.approx."))
2201 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2202 else if (Name.consume_front(
"atomic.load."))
2211 else if (Name.consume_front(
"atomic."))
2226 else if (Name.consume_front(
"bitcast."))
2229 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2230 else if (Name.consume_front(
"rotate."))
2232 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2233 else if (Name.consume_front(
"ptr.gen.to."))
2236 else if (Name.consume_front(
"ptr."))
2239 else if (Name.consume_front(
"ldg.global."))
2241 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2242 Name.starts_with(
"p."));
2245 .
Case(
"barrier0",
true)
2246 .
Case(
"barrier.n",
true)
2247 .
Case(
"barrier.sync.cnt",
true)
2248 .
Case(
"barrier.sync",
true)
2249 .
Case(
"barrier",
true)
2250 .
Case(
"bar.sync",
true)
2251 .
Case(
"barrier0.popc",
true)
2252 .
Case(
"barrier0.and",
true)
2253 .
Case(
"barrier0.or",
true)
2254 .
Case(
"clz.ll",
true)
2255 .
Case(
"popc.ll",
true)
2257 .
Case(
"swap.lo.hi.b64",
true)
2258 .
Case(
"tanh.approx.f32",
true)
2270 if (Name.starts_with(
"objectsize.")) {
2271 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2272 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2275 Intrinsic::objectsize, Tys);
2282 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2285 F->getParent(), Intrinsic::ptr_annotation,
2286 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2292 if (Name.consume_front(
"riscv.")) {
2295 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2296 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2297 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2298 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2301 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2314 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2315 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2324 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2325 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2326 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2327 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2332 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2341 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2343 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2352 if (Name ==
"stackprotectorcheck") {
2356 if (Name.starts_with(
"strip.invariant.group")) {
2361 F->getParent(), Intrinsic::launder_invariant_group,
2362 F->getReturnType());
2368 if (Name ==
"thread.pointer") {
2370 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2376 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2379 F->getParent(), Intrinsic::var_annotation,
2380 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2383 if (Name.consume_front(
"vector.splice")) {
2384 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2394 if (Name.consume_front(
"wasm.")) {
2397 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2398 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2399 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2404 F->getReturnType());
2408 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2410 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2412 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2431 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2441 std::string
Name =
F->getName().str();
2444 Name,
F->getParent());
2455 if (Result != std::nullopt) {
2472 bool CanUpgradeDebugIntrinsicsToRecords) {
2492 GV->
getName() ==
"llvm.global_dtors")) ||
2507 unsigned N =
Init->getNumOperands();
2508 std::vector<Constant *> NewCtors(
N);
2509 for (
unsigned i = 0; i !=
N; ++i) {
2512 Ctor->getAggregateElement(1),
2526 unsigned NumElts = ResultTy->getNumElements() * 8;
2530 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2540 for (
unsigned l = 0; l != NumElts; l += 16)
2541 for (
unsigned i = 0; i != 16; ++i) {
2542 unsigned Idx = NumElts + i - Shift;
2544 Idx -= NumElts - 16;
2545 Idxs[l + i] = Idx + l;
2548 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2552 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2560 unsigned NumElts = ResultTy->getNumElements() * 8;
2564 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2574 for (
unsigned l = 0; l != NumElts; l += 16)
2575 for (
unsigned i = 0; i != 16; ++i) {
2576 unsigned Idx = i + Shift;
2578 Idx += NumElts - 16;
2579 Idxs[l + i] = Idx + l;
2582 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2586 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2594 Mask = Builder.CreateBitCast(Mask, MaskTy);
2600 for (
unsigned i = 0; i != NumElts; ++i)
2602 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2613 if (
C->isAllOnesValue())
2618 return Builder.CreateSelect(Mask, Op0, Op1);
2625 if (
C->isAllOnesValue())
2629 Mask->getType()->getIntegerBitWidth());
2630 Mask = Builder.CreateBitCast(Mask, MaskTy);
2631 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2632 return Builder.CreateSelect(Mask, Op0, Op1);
2645 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2646 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2651 ShiftVal &= (NumElts - 1);
2660 if (ShiftVal > 16) {
2668 for (
unsigned l = 0; l < NumElts; l += 16) {
2669 for (
unsigned i = 0; i != 16; ++i) {
2670 unsigned Idx = ShiftVal + i;
2671 if (!IsVALIGN && Idx >= 16)
2672 Idx += NumElts - 16;
2673 Indices[l + i] = Idx + l;
2678 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2684 bool ZeroMask,
bool IndexForm) {
2687 unsigned EltWidth = Ty->getScalarSizeInBits();
2688 bool IsFloat = Ty->isFPOrFPVectorTy();
2690 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2691 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2692 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2693 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2694 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2695 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2696 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2697 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2698 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2699 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2700 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2701 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2702 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2703 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2704 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2705 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2706 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2707 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2708 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2709 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2710 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2711 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2712 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2713 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2714 else if (VecWidth == 128 && EltWidth == 16)
2715 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2716 else if (VecWidth == 256 && EltWidth == 16)
2717 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2718 else if (VecWidth == 512 && EltWidth == 16)
2719 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2720 else if (VecWidth == 128 && EltWidth == 8)
2721 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2722 else if (VecWidth == 256 && EltWidth == 8)
2723 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2724 else if (VecWidth == 512 && EltWidth == 8)
2725 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2736 Value *V = Builder.CreateIntrinsic(IID, Args);
2748 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2759 bool IsRotateRight) {
2769 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2770 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2773 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2774 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2819 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2824 bool IsShiftRight,
bool ZeroMask) {
2838 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2839 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2842 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2843 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2858 const Align Alignment =
2860 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2865 if (
C->isAllOnesValue())
2866 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2871 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2877 const Align Alignment =
2886 if (
C->isAllOnesValue())
2887 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2892 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2898 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2899 {Op0, Builder.getInt1(
false)});
2914 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2915 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2916 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2917 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2918 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2921 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2922 LHS = Builder.CreateAnd(
LHS, Mask);
2923 RHS = Builder.CreateAnd(
RHS, Mask);
2940 if (!
C || !
C->isAllOnesValue())
2941 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2946 for (
unsigned i = 0; i != NumElts; ++i)
2948 for (
unsigned i = NumElts; i != 8; ++i)
2949 Indices[i] = NumElts + i % NumElts;
2950 Vec = Builder.CreateShuffleVector(Vec,
2954 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2958 unsigned CC,
bool Signed) {
2966 }
else if (CC == 7) {
3002 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3003 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3005 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3006 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3015 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3021 Name = Name.substr(12);
3026 if (Name.starts_with(
"max.p")) {
3027 if (VecWidth == 128 && EltWidth == 32)
3028 IID = Intrinsic::x86_sse_max_ps;
3029 else if (VecWidth == 128 && EltWidth == 64)
3030 IID = Intrinsic::x86_sse2_max_pd;
3031 else if (VecWidth == 256 && EltWidth == 32)
3032 IID = Intrinsic::x86_avx_max_ps_256;
3033 else if (VecWidth == 256 && EltWidth == 64)
3034 IID = Intrinsic::x86_avx_max_pd_256;
3037 }
else if (Name.starts_with(
"min.p")) {
3038 if (VecWidth == 128 && EltWidth == 32)
3039 IID = Intrinsic::x86_sse_min_ps;
3040 else if (VecWidth == 128 && EltWidth == 64)
3041 IID = Intrinsic::x86_sse2_min_pd;
3042 else if (VecWidth == 256 && EltWidth == 32)
3043 IID = Intrinsic::x86_avx_min_ps_256;
3044 else if (VecWidth == 256 && EltWidth == 64)
3045 IID = Intrinsic::x86_avx_min_pd_256;
3048 }
else if (Name.starts_with(
"pshuf.b.")) {
3049 if (VecWidth == 128)
3050 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3051 else if (VecWidth == 256)
3052 IID = Intrinsic::x86_avx2_pshuf_b;
3053 else if (VecWidth == 512)
3054 IID = Intrinsic::x86_avx512_pshuf_b_512;
3057 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3058 if (VecWidth == 128)
3059 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3060 else if (VecWidth == 256)
3061 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3062 else if (VecWidth == 512)
3063 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3066 }
else if (Name.starts_with(
"pmulh.w")) {
3067 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3068 "Unexpected intrinsic");
3071 }
else if (Name.starts_with(
"pmulhu.w")) {
3072 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3073 "Unexpected intrinsic");
3076 }
else if (Name.starts_with(
"pmaddw.d.")) {
3077 if (VecWidth == 128)
3078 IID = Intrinsic::x86_sse2_pmadd_wd;
3079 else if (VecWidth == 256)
3080 IID = Intrinsic::x86_avx2_pmadd_wd;
3081 else if (VecWidth == 512)
3082 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3085 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3086 if (VecWidth == 128)
3087 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3088 else if (VecWidth == 256)
3089 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3090 else if (VecWidth == 512)
3091 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3094 }
else if (Name.starts_with(
"packsswb.")) {
3095 if (VecWidth == 128)
3096 IID = Intrinsic::x86_sse2_packsswb_128;
3097 else if (VecWidth == 256)
3098 IID = Intrinsic::x86_avx2_packsswb;
3099 else if (VecWidth == 512)
3100 IID = Intrinsic::x86_avx512_packsswb_512;
3103 }
else if (Name.starts_with(
"packssdw.")) {
3104 if (VecWidth == 128)
3105 IID = Intrinsic::x86_sse2_packssdw_128;
3106 else if (VecWidth == 256)
3107 IID = Intrinsic::x86_avx2_packssdw;
3108 else if (VecWidth == 512)
3109 IID = Intrinsic::x86_avx512_packssdw_512;
3112 }
else if (Name.starts_with(
"packuswb.")) {
3113 if (VecWidth == 128)
3114 IID = Intrinsic::x86_sse2_packuswb_128;
3115 else if (VecWidth == 256)
3116 IID = Intrinsic::x86_avx2_packuswb;
3117 else if (VecWidth == 512)
3118 IID = Intrinsic::x86_avx512_packuswb_512;
3121 }
else if (Name.starts_with(
"packusdw.")) {
3122 if (VecWidth == 128)
3123 IID = Intrinsic::x86_sse41_packusdw;
3124 else if (VecWidth == 256)
3125 IID = Intrinsic::x86_avx2_packusdw;
3126 else if (VecWidth == 512)
3127 IID = Intrinsic::x86_avx512_packusdw_512;
3130 }
else if (Name.starts_with(
"vpermilvar.")) {
3131 if (VecWidth == 128 && EltWidth == 32)
3132 IID = Intrinsic::x86_avx_vpermilvar_ps;
3133 else if (VecWidth == 128 && EltWidth == 64)
3134 IID = Intrinsic::x86_avx_vpermilvar_pd;
3135 else if (VecWidth == 256 && EltWidth == 32)
3136 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3137 else if (VecWidth == 256 && EltWidth == 64)
3138 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3139 else if (VecWidth == 512 && EltWidth == 32)
3140 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3141 else if (VecWidth == 512 && EltWidth == 64)
3142 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3145 }
else if (Name ==
"cvtpd2dq.256") {
3146 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3147 }
else if (Name ==
"cvtpd2ps.256") {
3148 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3149 }
else if (Name ==
"cvttpd2dq.256") {
3150 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3151 }
else if (Name ==
"cvttps2dq.128") {
3152 IID = Intrinsic::x86_sse2_cvttps2dq;
3153 }
else if (Name ==
"cvttps2dq.256") {
3154 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3155 }
else if (Name.starts_with(
"permvar.")) {
3157 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3158 IID = Intrinsic::x86_avx2_permps;
3159 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3160 IID = Intrinsic::x86_avx2_permd;
3161 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3162 IID = Intrinsic::x86_avx512_permvar_df_256;
3163 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3164 IID = Intrinsic::x86_avx512_permvar_di_256;
3165 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3166 IID = Intrinsic::x86_avx512_permvar_sf_512;
3167 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3168 IID = Intrinsic::x86_avx512_permvar_si_512;
3169 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3170 IID = Intrinsic::x86_avx512_permvar_df_512;
3171 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3172 IID = Intrinsic::x86_avx512_permvar_di_512;
3173 else if (VecWidth == 128 && EltWidth == 16)
3174 IID = Intrinsic::x86_avx512_permvar_hi_128;
3175 else if (VecWidth == 256 && EltWidth == 16)
3176 IID = Intrinsic::x86_avx512_permvar_hi_256;
3177 else if (VecWidth == 512 && EltWidth == 16)
3178 IID = Intrinsic::x86_avx512_permvar_hi_512;
3179 else if (VecWidth == 128 && EltWidth == 8)
3180 IID = Intrinsic::x86_avx512_permvar_qi_128;
3181 else if (VecWidth == 256 && EltWidth == 8)
3182 IID = Intrinsic::x86_avx512_permvar_qi_256;
3183 else if (VecWidth == 512 && EltWidth == 8)
3184 IID = Intrinsic::x86_avx512_permvar_qi_512;
3187 }
else if (Name.starts_with(
"dbpsadbw.")) {
3188 if (VecWidth == 128)
3189 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3190 else if (VecWidth == 256)
3191 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3192 else if (VecWidth == 512)
3193 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3196 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3197 if (VecWidth == 128)
3198 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3199 else if (VecWidth == 256)
3200 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3201 else if (VecWidth == 512)
3202 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3205 }
else if (Name.starts_with(
"conflict.")) {
3206 if (Name[9] ==
'd' && VecWidth == 128)
3207 IID = Intrinsic::x86_avx512_conflict_d_128;
3208 else if (Name[9] ==
'd' && VecWidth == 256)
3209 IID = Intrinsic::x86_avx512_conflict_d_256;
3210 else if (Name[9] ==
'd' && VecWidth == 512)
3211 IID = Intrinsic::x86_avx512_conflict_d_512;
3212 else if (Name[9] ==
'q' && VecWidth == 128)
3213 IID = Intrinsic::x86_avx512_conflict_q_128;
3214 else if (Name[9] ==
'q' && VecWidth == 256)
3215 IID = Intrinsic::x86_avx512_conflict_q_256;
3216 else if (Name[9] ==
'q' && VecWidth == 512)
3217 IID = Intrinsic::x86_avx512_conflict_q_512;
3220 }
else if (Name.starts_with(
"pavg.")) {
3221 if (Name[5] ==
'b' && VecWidth == 128)
3222 IID = Intrinsic::x86_sse2_pavg_b;
3223 else if (Name[5] ==
'b' && VecWidth == 256)
3224 IID = Intrinsic::x86_avx2_pavg_b;
3225 else if (Name[5] ==
'b' && VecWidth == 512)
3226 IID = Intrinsic::x86_avx512_pavg_b_512;
3227 else if (Name[5] ==
'w' && VecWidth == 128)
3228 IID = Intrinsic::x86_sse2_pavg_w;
3229 else if (Name[5] ==
'w' && VecWidth == 256)
3230 IID = Intrinsic::x86_avx2_pavg_w;
3231 else if (Name[5] ==
'w' && VecWidth == 512)
3232 IID = Intrinsic::x86_avx512_pavg_w_512;
3241 Rep = Builder.CreateIntrinsic(IID, Args);
3252 if (AsmStr->find(
"mov\tfp") == 0 &&
3253 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3254 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3255 AsmStr->replace(Pos, 1,
";");
3261 Value *Rep =
nullptr;
3263 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3265 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3266 {Arg, Builder.getTrue()},
3268 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3269 Type *Ty = (Name ==
"abs.bf16")
3273 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3274 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3275 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3276 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3277 : Intrinsic::nvvm_fabs;
3278 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3279 }
else if (Name.consume_front(
"add.")) {
3282 assert(
FAdd &&
"unsupported nvvm.add.* intrinsic");
3285 Rep = Builder.CreateIntrinsic(
3287 {A, CI->getArgOperand(1),
3288 Builder.getInt32(static_cast<int>(RoundingMode))});
3289 }
else if (Name.consume_front(
"ex2.approx.")) {
3291 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3292 : Intrinsic::nvvm_ex2_approx;
3293 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3294 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3295 Name.starts_with(
"atomic.load.add.f64.p")) {
3298 Rep = Builder.CreateAtomicRMW(
3304 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3305 Name.starts_with(
"atomic.load.dec.32.p")) {
3310 Rep = Builder.CreateAtomicRMW(
3314 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3320 Op.contains(
".cta.") ?
"block" :
"");
3321 if (
Op.starts_with(
"cas.")) {
3323 Value *Pair = Builder.CreateAtomicCmpXchg(
3326 Rep = Builder.CreateExtractValue(Pair, 0);
3344 "unexpected nvvm scoped atomic intrinsic");
3345 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3348 }
else if (Name ==
"clz.ll") {
3351 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3352 {Arg, Builder.getFalse()},
3354 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3355 }
else if (Name ==
"popc.ll") {
3359 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3360 Arg,
nullptr,
"ctpop");
3361 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3362 }
else if (Name ==
"h2f") {
3364 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3365 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3366 }
else if (Name.consume_front(
"bitcast.") &&
3367 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3370 }
else if (Name ==
"rotate.b32") {
3373 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3374 {Arg, Arg, ShiftAmt});
3375 }
else if (Name ==
"rotate.b64") {
3379 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3380 {Arg, Arg, ZExtShiftAmt});
3381 }
else if (Name ==
"rotate.right.b64") {
3385 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3386 {Arg, Arg, ZExtShiftAmt});
3387 }
else if (Name ==
"swap.lo.hi.b64") {
3390 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3391 {Arg, Arg, Builder.getInt64(32)});
3392 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3395 Name.starts_with(
".to.gen"))) {
3397 }
else if (Name.consume_front(
"ldg.global")) {
3401 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3404 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3406 }
else if (Name ==
"tanh.approx.f32") {
3410 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3412 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3414 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3415 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3417 }
else if (Name ==
"barrier") {
3418 Rep = Builder.CreateIntrinsic(
3419 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3421 }
else if (Name ==
"barrier.sync") {
3422 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3424 }
else if (Name ==
"barrier.sync.cnt") {
3425 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3427 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3428 Name ==
"barrier0.or") {
3430 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3434 .
Case(
"barrier0.popc",
3435 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3436 .
Case(
"barrier0.and",
3437 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3438 .
Case(
"barrier0.or",
3439 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3440 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3441 Rep = Builder.CreateZExt(Bar, CI->
getType());
3455 ? Builder.CreateBitCast(Arg, NewType)
3458 Rep = Builder.CreateCall(NewFn, Args);
3459 if (
F->getReturnType()->isIntegerTy())
3460 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3470 Value *Rep =
nullptr;
3472 if (Name.starts_with(
"sse4a.movnt.")) {
3484 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3487 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3488 }
else if (Name.starts_with(
"avx.movnt.") ||
3489 Name.starts_with(
"avx512.storent.")) {
3501 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3502 }
else if (Name ==
"sse2.storel.dq") {
3507 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3508 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3509 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3510 }
else if (Name.starts_with(
"sse.storeu.") ||
3511 Name.starts_with(
"sse2.storeu.") ||
3512 Name.starts_with(
"avx.storeu.")) {
3515 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3516 }
else if (Name ==
"avx512.mask.store.ss") {
3520 }
else if (Name.starts_with(
"avx512.mask.store")) {
3522 bool Aligned = Name[17] !=
'u';
3525 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3528 bool CmpEq = Name[9] ==
'e';
3531 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3532 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3539 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3540 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3542 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3543 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3544 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3545 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3546 Name.starts_with(
"sse2.sqrt.p") ||
3547 Name.starts_with(
"sse.sqrt.p")) {
3548 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3549 {CI->getArgOperand(0)});
3550 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3554 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3555 : Intrinsic::x86_avx512_sqrt_pd_512;
3558 Rep = Builder.CreateIntrinsic(IID, Args);
3560 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3561 {CI->getArgOperand(0)});
3565 }
else if (Name.starts_with(
"avx512.ptestm") ||
3566 Name.starts_with(
"avx512.ptestnm")) {
3570 Rep = Builder.CreateAnd(Op0, Op1);
3576 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3578 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3581 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3584 }
else if (Name.starts_with(
"avx512.kunpck")) {
3589 for (
unsigned i = 0; i != NumElts; ++i)
3598 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3599 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3600 }
else if (Name ==
"avx512.kand.w") {
3603 Rep = Builder.CreateAnd(
LHS,
RHS);
3604 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3605 }
else if (Name ==
"avx512.kandn.w") {
3608 LHS = Builder.CreateNot(
LHS);
3609 Rep = Builder.CreateAnd(
LHS,
RHS);
3610 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3611 }
else if (Name ==
"avx512.kor.w") {
3614 Rep = Builder.CreateOr(
LHS,
RHS);
3615 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3616 }
else if (Name ==
"avx512.kxor.w") {
3619 Rep = Builder.CreateXor(
LHS,
RHS);
3620 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3621 }
else if (Name ==
"avx512.kxnor.w") {
3624 LHS = Builder.CreateNot(
LHS);
3625 Rep = Builder.CreateXor(
LHS,
RHS);
3626 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3627 }
else if (Name ==
"avx512.knot.w") {
3629 Rep = Builder.CreateNot(Rep);
3630 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3631 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3634 Rep = Builder.CreateOr(
LHS,
RHS);
3635 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3637 if (Name[14] ==
'c')
3641 Rep = Builder.CreateICmpEQ(Rep,
C);
3642 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3643 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3644 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3645 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3646 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3649 ConstantInt::get(I32Ty, 0));
3651 ConstantInt::get(I32Ty, 0));
3653 if (Name.contains(
".add."))
3654 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3655 else if (Name.contains(
".sub."))
3656 EltOp = Builder.CreateFSub(Elt0, Elt1);
3657 else if (Name.contains(
".mul."))
3658 EltOp = Builder.CreateFMul(Elt0, Elt1);
3660 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3661 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3662 ConstantInt::get(I32Ty, 0));
3663 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3665 bool CmpEq = Name[16] ==
'e';
3667 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3669 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3676 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3679 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3682 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3689 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3691 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3692 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3694 if (VecWidth == 128 && EltWidth == 32)
3695 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3696 else if (VecWidth == 256 && EltWidth == 32)
3697 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3698 else if (VecWidth == 512 && EltWidth == 32)
3699 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3700 else if (VecWidth == 128 && EltWidth == 64)
3701 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3702 else if (VecWidth == 256 && EltWidth == 64)
3703 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3704 else if (VecWidth == 512 && EltWidth == 64)
3705 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3712 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3715 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3716 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3718 if (VecWidth == 128 && EltWidth == 32)
3719 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3720 else if (VecWidth == 256 && EltWidth == 32)
3721 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3722 else if (VecWidth == 512 && EltWidth == 32)
3723 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3724 else if (VecWidth == 128 && EltWidth == 64)
3725 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3726 else if (VecWidth == 256 && EltWidth == 64)
3727 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3728 else if (VecWidth == 512 && EltWidth == 64)
3729 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3734 if (VecWidth == 512)
3736 Args.push_back(Mask);
3738 Rep = Builder.CreateIntrinsic(IID, Args);
3739 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3743 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3746 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3747 Name.starts_with(
"avx512.cvtw2mask.") ||
3748 Name.starts_with(
"avx512.cvtd2mask.") ||
3749 Name.starts_with(
"avx512.cvtq2mask.")) {
3754 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3755 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3756 Name.starts_with(
"avx512.mask.pabs")) {
3758 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3759 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3760 Name.starts_with(
"avx512.mask.pmaxs")) {
3762 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3763 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3764 Name.starts_with(
"avx512.mask.pmaxu")) {
3766 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3767 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3768 Name.starts_with(
"avx512.mask.pmins")) {
3770 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3771 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3772 Name.starts_with(
"avx512.mask.pminu")) {
3774 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3775 Name.starts_with(
"avx512.pmulh.w")) {
3777 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3778 Name.starts_with(
"avx512.pmulhu.w")) {
3780 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3781 Name ==
"avx512.pmulu.dq.512" ||
3782 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3784 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3785 Name ==
"avx512.pmul.dq.512" ||
3786 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3788 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3789 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3794 }
else if (Name ==
"avx512.cvtusi2sd") {
3799 }
else if (Name ==
"sse2.cvtss2sd") {
3801 Rep = Builder.CreateFPExt(
3804 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3805 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3806 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3807 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3808 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3809 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3810 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3811 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3812 Name ==
"avx512.mask.cvtqq2ps.256" ||
3813 Name ==
"avx512.mask.cvtqq2ps.512" ||
3814 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3815 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3816 Name ==
"avx.cvt.ps2.pd.256" ||
3817 Name ==
"avx512.mask.cvtps2pd.128" ||
3818 Name ==
"avx512.mask.cvtps2pd.256") {
3823 unsigned NumDstElts = DstTy->getNumElements();
3824 if (NumDstElts < SrcTy->getNumElements()) {
3825 assert(NumDstElts == 2 &&
"Unexpected vector size");
3826 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3829 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3830 bool IsUnsigned = Name.contains(
"cvtu");
3832 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3836 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3837 : Intrinsic::x86_avx512_sitofp_round;
3838 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3841 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3842 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3848 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3849 Name.starts_with(
"vcvtph2ps.")) {
3853 unsigned NumDstElts = DstTy->getNumElements();
3854 if (NumDstElts != SrcTy->getNumElements()) {
3855 assert(NumDstElts == 4 &&
"Unexpected vector size");
3856 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3858 Rep = Builder.CreateBitCast(
3860 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3864 }
else if (Name.starts_with(
"avx512.mask.load")) {
3866 bool Aligned = Name[16] !=
'u';
3869 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3873 ResultTy->getNumElements());
3874 Rep = Builder.CreateIntrinsic(
3875 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3877 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3883 Rep = Builder.CreateIntrinsic(
3884 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3886 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3887 Name.starts_with(
"avx512.mask.expand.")) {
3891 ResultTy->getNumElements());
3893 bool IsCompress = Name[12] ==
'c';
3894 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3895 : Intrinsic::x86_avx512_mask_expand;
3896 Rep = Builder.CreateIntrinsic(
3898 }
else if (Name.starts_with(
"xop.vpcom")) {
3900 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3901 Name.ends_with(
"uq"))
3903 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3904 Name.ends_with(
"d") || Name.ends_with(
"q"))
3913 Name = Name.substr(9);
3914 if (Name.starts_with(
"lt"))
3916 else if (Name.starts_with(
"le"))
3918 else if (Name.starts_with(
"gt"))
3920 else if (Name.starts_with(
"ge"))
3922 else if (Name.starts_with(
"eq"))
3924 else if (Name.starts_with(
"ne"))
3926 else if (Name.starts_with(
"false"))
3928 else if (Name.starts_with(
"true"))
3935 }
else if (Name.starts_with(
"xop.vpcmov")) {
3937 Value *NotSel = Builder.CreateNot(Sel);
3940 Rep = Builder.CreateOr(Sel0, Sel1);
3941 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3942 Name.starts_with(
"avx512.mask.prol")) {
3944 }
else if (Name.starts_with(
"avx512.pror") ||
3945 Name.starts_with(
"avx512.mask.pror")) {
3947 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3948 Name.starts_with(
"avx512.mask.vpshld") ||
3949 Name.starts_with(
"avx512.maskz.vpshld")) {
3950 bool ZeroMask = Name[11] ==
'z';
3952 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3953 Name.starts_with(
"avx512.mask.vpshrd") ||
3954 Name.starts_with(
"avx512.maskz.vpshrd")) {
3955 bool ZeroMask = Name[11] ==
'z';
3957 }
else if (Name ==
"sse42.crc32.64.8") {
3960 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3962 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3963 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3964 Name.starts_with(
"avx512.vbroadcast.s")) {
3967 Type *EltTy = VecTy->getElementType();
3968 unsigned EltNum = VecTy->getNumElements();
3972 for (
unsigned I = 0;
I < EltNum; ++
I)
3973 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3974 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3975 Name.starts_with(
"sse41.pmovzx") ||
3976 Name.starts_with(
"avx2.pmovsx") ||
3977 Name.starts_with(
"avx2.pmovzx") ||
3978 Name.starts_with(
"avx512.mask.pmovsx") ||
3979 Name.starts_with(
"avx512.mask.pmovzx")) {
3981 unsigned NumDstElts = DstTy->getNumElements();
3985 for (
unsigned i = 0; i != NumDstElts; ++i)
3990 bool DoSext = Name.contains(
"pmovsx");
3992 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3997 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3998 Name ==
"avx512.mask.pmov.qd.512" ||
3999 Name ==
"avx512.mask.pmov.wb.256" ||
4000 Name ==
"avx512.mask.pmov.wb.512") {
4005 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4006 Name ==
"avx2.vbroadcasti128") {
4012 if (NumSrcElts == 2)
4015 Rep = Builder.CreateShuffleVector(
Load,
4017 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4018 Name.starts_with(
"avx512.mask.shuf.f")) {
4023 unsigned ControlBitsMask = NumLanes - 1;
4024 unsigned NumControlBits = NumLanes / 2;
4027 for (
unsigned l = 0; l != NumLanes; ++l) {
4028 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4030 if (l >= NumLanes / 2)
4031 LaneMask += NumLanes;
4032 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4033 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4039 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4040 Name.starts_with(
"avx512.mask.broadcasti")) {
4043 unsigned NumDstElts =
4047 for (
unsigned i = 0; i != NumDstElts; ++i)
4048 ShuffleMask[i] = i % NumSrcElts;
4054 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4055 Name.starts_with(
"avx2.vbroadcast") ||
4056 Name.starts_with(
"avx512.pbroadcast") ||
4057 Name.starts_with(
"avx512.mask.broadcast.s")) {
4064 Rep = Builder.CreateShuffleVector(
Op, M);
4069 }
else if (Name.starts_with(
"sse2.padds.") ||
4070 Name.starts_with(
"avx2.padds.") ||
4071 Name.starts_with(
"avx512.padds.") ||
4072 Name.starts_with(
"avx512.mask.padds.")) {
4074 }
else if (Name.starts_with(
"sse2.psubs.") ||
4075 Name.starts_with(
"avx2.psubs.") ||
4076 Name.starts_with(
"avx512.psubs.") ||
4077 Name.starts_with(
"avx512.mask.psubs.")) {
4079 }
else if (Name.starts_with(
"sse2.paddus.") ||
4080 Name.starts_with(
"avx2.paddus.") ||
4081 Name.starts_with(
"avx512.mask.paddus.")) {
4083 }
else if (Name.starts_with(
"sse2.psubus.") ||
4084 Name.starts_with(
"avx2.psubus.") ||
4085 Name.starts_with(
"avx512.mask.psubus.")) {
4087 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4092 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4096 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4101 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4106 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4107 Name ==
"avx512.psll.dq.512") {
4111 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4112 Name ==
"avx512.psrl.dq.512") {
4116 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4117 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4118 Name.starts_with(
"avx2.pblendd.")) {
4123 unsigned NumElts = VecTy->getNumElements();
4126 for (
unsigned i = 0; i != NumElts; ++i)
4127 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4129 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4130 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4131 Name ==
"avx2.vinserti128" ||
4132 Name.starts_with(
"avx512.mask.insert")) {
4136 unsigned DstNumElts =
4138 unsigned SrcNumElts =
4140 unsigned Scale = DstNumElts / SrcNumElts;
4147 for (
unsigned i = 0; i != SrcNumElts; ++i)
4149 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4150 Idxs[i] = SrcNumElts;
4151 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4165 for (
unsigned i = 0; i != DstNumElts; ++i)
4168 for (
unsigned i = 0; i != SrcNumElts; ++i)
4169 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4170 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4176 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4177 Name ==
"avx2.vextracti128" ||
4178 Name.starts_with(
"avx512.mask.vextract")) {
4181 unsigned DstNumElts =
4183 unsigned SrcNumElts =
4185 unsigned Scale = SrcNumElts / DstNumElts;
4192 for (
unsigned i = 0; i != DstNumElts; ++i) {
4193 Idxs[i] = i + (
Imm * DstNumElts);
4195 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4201 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4202 Name.starts_with(
"avx512.mask.perm.di.")) {
4206 unsigned NumElts = VecTy->getNumElements();
4209 for (
unsigned i = 0; i != NumElts; ++i)
4210 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4212 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4217 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4229 unsigned HalfSize = NumElts / 2;
4241 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4242 for (
unsigned i = 0; i < HalfSize; ++i)
4243 ShuffleMask[i] = StartIndex + i;
4246 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4247 for (
unsigned i = 0; i < HalfSize; ++i)
4248 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4250 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4252 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4253 Name.starts_with(
"avx512.mask.vpermil.p") ||
4254 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4258 unsigned NumElts = VecTy->getNumElements();
4260 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4261 unsigned IdxMask = ((1 << IdxSize) - 1);
4267 for (
unsigned i = 0; i != NumElts; ++i)
4268 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4270 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4275 }
else if (Name ==
"sse2.pshufl.w" ||
4276 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4281 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4285 for (
unsigned l = 0; l != NumElts; l += 8) {
4286 for (
unsigned i = 0; i != 4; ++i)
4287 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4288 for (
unsigned i = 4; i != 8; ++i)
4289 Idxs[i + l] = i + l;
4292 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4297 }
else if (Name ==
"sse2.pshufh.w" ||
4298 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4303 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4307 for (
unsigned l = 0; l != NumElts; l += 8) {
4308 for (
unsigned i = 0; i != 4; ++i)
4309 Idxs[i + l] = i + l;
4310 for (
unsigned i = 0; i != 4; ++i)
4311 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4314 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4319 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4326 unsigned HalfLaneElts = NumLaneElts / 2;
4329 for (
unsigned i = 0; i != NumElts; ++i) {
4331 Idxs[i] = i - (i % NumLaneElts);
4333 if ((i % NumLaneElts) >= HalfLaneElts)
4337 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4340 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4344 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4345 Name.starts_with(
"avx512.mask.movshdup") ||
4346 Name.starts_with(
"avx512.mask.movsldup")) {
4352 if (Name.starts_with(
"avx512.mask.movshdup."))
4356 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4357 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4358 Idxs[i + l + 0] = i + l +
Offset;
4359 Idxs[i + l + 1] = i + l +
Offset;
4362 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4366 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4367 Name.starts_with(
"avx512.mask.unpckl.")) {
4374 for (
int l = 0; l != NumElts; l += NumLaneElts)
4375 for (
int i = 0; i != NumLaneElts; ++i)
4376 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4378 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4382 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4383 Name.starts_with(
"avx512.mask.unpckh.")) {
4390 for (
int l = 0; l != NumElts; l += NumLaneElts)
4391 for (
int i = 0; i != NumLaneElts; ++i)
4392 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4394 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4398 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4399 Name.starts_with(
"avx512.mask.pand.")) {
4402 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4404 Rep = Builder.CreateBitCast(Rep, FTy);
4407 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4408 Name.starts_with(
"avx512.mask.pandn.")) {
4411 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4412 Rep = Builder.CreateAnd(Rep,
4414 Rep = Builder.CreateBitCast(Rep, FTy);
4417 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4418 Name.starts_with(
"avx512.mask.por.")) {
4421 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4423 Rep = Builder.CreateBitCast(Rep, FTy);
4426 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4427 Name.starts_with(
"avx512.mask.pxor.")) {
4430 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4432 Rep = Builder.CreateBitCast(Rep, FTy);
4435 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4439 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4443 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4447 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4448 if (Name.ends_with(
".512")) {
4450 if (Name[17] ==
's')
4451 IID = Intrinsic::x86_avx512_add_ps_512;
4453 IID = Intrinsic::x86_avx512_add_pd_512;
4455 Rep = Builder.CreateIntrinsic(
4463 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4464 if (Name.ends_with(
".512")) {
4466 if (Name[17] ==
's')
4467 IID = Intrinsic::x86_avx512_div_ps_512;
4469 IID = Intrinsic::x86_avx512_div_pd_512;
4471 Rep = Builder.CreateIntrinsic(
4479 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4480 if (Name.ends_with(
".512")) {
4482 if (Name[17] ==
's')
4483 IID = Intrinsic::x86_avx512_mul_ps_512;
4485 IID = Intrinsic::x86_avx512_mul_pd_512;
4487 Rep = Builder.CreateIntrinsic(
4495 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4496 if (Name.ends_with(
".512")) {
4498 if (Name[17] ==
's')
4499 IID = Intrinsic::x86_avx512_sub_ps_512;
4501 IID = Intrinsic::x86_avx512_sub_pd_512;
4503 Rep = Builder.CreateIntrinsic(
4511 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4512 Name.starts_with(
"avx512.mask.min.p")) &&
4513 Name.drop_front(18) ==
".512") {
4514 bool IsDouble = Name[17] ==
'd';
4515 bool IsMin = Name[13] ==
'i';
4517 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4518 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4521 Rep = Builder.CreateIntrinsic(
4526 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4528 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4529 {CI->getArgOperand(0), Builder.getInt1(false)});
4532 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4533 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4534 bool IsVariable = Name[16] ==
'v';
4535 char Size = Name[16] ==
'.' ? Name[17]
4536 : Name[17] ==
'.' ? Name[18]
4537 : Name[18] ==
'.' ? Name[19]
4541 if (IsVariable && Name[17] !=
'.') {
4542 if (
Size ==
'd' && Name[17] ==
'2')
4543 IID = Intrinsic::x86_avx2_psllv_q;
4544 else if (
Size ==
'd' && Name[17] ==
'4')
4545 IID = Intrinsic::x86_avx2_psllv_q_256;
4546 else if (
Size ==
's' && Name[17] ==
'4')
4547 IID = Intrinsic::x86_avx2_psllv_d;
4548 else if (
Size ==
's' && Name[17] ==
'8')
4549 IID = Intrinsic::x86_avx2_psllv_d_256;
4550 else if (
Size ==
'h' && Name[17] ==
'8')
4551 IID = Intrinsic::x86_avx512_psllv_w_128;
4552 else if (
Size ==
'h' && Name[17] ==
'1')
4553 IID = Intrinsic::x86_avx512_psllv_w_256;
4554 else if (Name[17] ==
'3' && Name[18] ==
'2')
4555 IID = Intrinsic::x86_avx512_psllv_w_512;
4558 }
else if (Name.ends_with(
".128")) {
4560 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4561 : Intrinsic::x86_sse2_psll_d;
4562 else if (
Size ==
'q')
4563 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4564 : Intrinsic::x86_sse2_psll_q;
4565 else if (
Size ==
'w')
4566 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4567 : Intrinsic::x86_sse2_psll_w;
4570 }
else if (Name.ends_with(
".256")) {
4572 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4573 : Intrinsic::x86_avx2_psll_d;
4574 else if (
Size ==
'q')
4575 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4576 : Intrinsic::x86_avx2_psll_q;
4577 else if (
Size ==
'w')
4578 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4579 : Intrinsic::x86_avx2_psll_w;
4584 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4585 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4586 : Intrinsic::x86_avx512_psll_d_512;
4587 else if (
Size ==
'q')
4588 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4589 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4590 : Intrinsic::x86_avx512_psll_q_512;
4591 else if (
Size ==
'w')
4592 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4593 : Intrinsic::x86_avx512_psll_w_512;
4599 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4600 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4601 bool IsVariable = Name[16] ==
'v';
4602 char Size = Name[16] ==
'.' ? Name[17]
4603 : Name[17] ==
'.' ? Name[18]
4604 : Name[18] ==
'.' ? Name[19]
4608 if (IsVariable && Name[17] !=
'.') {
4609 if (
Size ==
'd' && Name[17] ==
'2')
4610 IID = Intrinsic::x86_avx2_psrlv_q;
4611 else if (
Size ==
'd' && Name[17] ==
'4')
4612 IID = Intrinsic::x86_avx2_psrlv_q_256;
4613 else if (
Size ==
's' && Name[17] ==
'4')
4614 IID = Intrinsic::x86_avx2_psrlv_d;
4615 else if (
Size ==
's' && Name[17] ==
'8')
4616 IID = Intrinsic::x86_avx2_psrlv_d_256;
4617 else if (
Size ==
'h' && Name[17] ==
'8')
4618 IID = Intrinsic::x86_avx512_psrlv_w_128;
4619 else if (
Size ==
'h' && Name[17] ==
'1')
4620 IID = Intrinsic::x86_avx512_psrlv_w_256;
4621 else if (Name[17] ==
'3' && Name[18] ==
'2')
4622 IID = Intrinsic::x86_avx512_psrlv_w_512;
4625 }
else if (Name.ends_with(
".128")) {
4627 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4628 : Intrinsic::x86_sse2_psrl_d;
4629 else if (
Size ==
'q')
4630 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4631 : Intrinsic::x86_sse2_psrl_q;
4632 else if (
Size ==
'w')
4633 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4634 : Intrinsic::x86_sse2_psrl_w;
4637 }
else if (Name.ends_with(
".256")) {
4639 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4640 : Intrinsic::x86_avx2_psrl_d;
4641 else if (
Size ==
'q')
4642 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4643 : Intrinsic::x86_avx2_psrl_q;
4644 else if (
Size ==
'w')
4645 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4646 : Intrinsic::x86_avx2_psrl_w;
4651 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4652 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4653 : Intrinsic::x86_avx512_psrl_d_512;
4654 else if (
Size ==
'q')
4655 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4656 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4657 : Intrinsic::x86_avx512_psrl_q_512;
4658 else if (
Size ==
'w')
4659 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4660 : Intrinsic::x86_avx512_psrl_w_512;
4666 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4667 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4668 bool IsVariable = Name[16] ==
'v';
4669 char Size = Name[16] ==
'.' ? Name[17]
4670 : Name[17] ==
'.' ? Name[18]
4671 : Name[18] ==
'.' ? Name[19]
4675 if (IsVariable && Name[17] !=
'.') {
4676 if (
Size ==
's' && Name[17] ==
'4')
4677 IID = Intrinsic::x86_avx2_psrav_d;
4678 else if (
Size ==
's' && Name[17] ==
'8')
4679 IID = Intrinsic::x86_avx2_psrav_d_256;
4680 else if (
Size ==
'h' && Name[17] ==
'8')
4681 IID = Intrinsic::x86_avx512_psrav_w_128;
4682 else if (
Size ==
'h' && Name[17] ==
'1')
4683 IID = Intrinsic::x86_avx512_psrav_w_256;
4684 else if (Name[17] ==
'3' && Name[18] ==
'2')
4685 IID = Intrinsic::x86_avx512_psrav_w_512;
4688 }
else if (Name.ends_with(
".128")) {
4690 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4691 : Intrinsic::x86_sse2_psra_d;
4692 else if (
Size ==
'q')
4693 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4694 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4695 : Intrinsic::x86_avx512_psra_q_128;
4696 else if (
Size ==
'w')
4697 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4698 : Intrinsic::x86_sse2_psra_w;
4701 }
else if (Name.ends_with(
".256")) {
4703 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4704 : Intrinsic::x86_avx2_psra_d;
4705 else if (
Size ==
'q')
4706 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4707 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4708 : Intrinsic::x86_avx512_psra_q_256;
4709 else if (
Size ==
'w')
4710 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4711 : Intrinsic::x86_avx2_psra_w;
4716 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4717 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4718 : Intrinsic::x86_avx512_psra_d_512;
4719 else if (
Size ==
'q')
4720 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4721 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4722 : Intrinsic::x86_avx512_psra_q_512;
4723 else if (
Size ==
'w')
4724 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4725 : Intrinsic::x86_avx512_psra_w_512;
4731 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4733 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4735 }
else if (Name.ends_with(
".movntdqa")) {
4739 LoadInst *LI = Builder.CreateAlignedLoad(
4744 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4745 Name.starts_with(
"fma.vfmsub.") ||
4746 Name.starts_with(
"fma.vfnmadd.") ||
4747 Name.starts_with(
"fma.vfnmsub.")) {
4748 bool NegMul = Name[6] ==
'n';
4749 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4750 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4761 if (NegMul && !IsScalar)
4762 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4763 if (NegMul && IsScalar)
4764 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4766 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4768 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4772 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4780 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4784 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4785 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4786 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4787 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4788 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4789 bool IsMask3 = Name[11] ==
'3';
4790 bool IsMaskZ = Name[11] ==
'z';
4792 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4793 bool NegMul = Name[2] ==
'n';
4794 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4800 if (NegMul && (IsMask3 || IsMaskZ))
4801 A = Builder.CreateFNeg(
A);
4802 if (NegMul && !(IsMask3 || IsMaskZ))
4803 B = Builder.CreateFNeg(
B);
4805 C = Builder.CreateFNeg(
C);
4807 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4808 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4809 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4816 if (Name.back() ==
'd')
4817 IID = Intrinsic::x86_avx512_vfmadd_f64;
4819 IID = Intrinsic::x86_avx512_vfmadd_f32;
4820 Rep = Builder.CreateIntrinsic(IID,
Ops);
4822 Rep = Builder.CreateFMA(
A,
B,
C);
4831 if (NegAcc && IsMask3)
4836 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4838 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4839 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4840 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4841 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4842 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4843 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4844 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4845 bool IsMask3 = Name[11] ==
'3';
4846 bool IsMaskZ = Name[11] ==
'z';
4848 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4849 bool NegMul = Name[2] ==
'n';
4850 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4856 if (NegMul && (IsMask3 || IsMaskZ))
4857 A = Builder.CreateFNeg(
A);
4858 if (NegMul && !(IsMask3 || IsMaskZ))
4859 B = Builder.CreateFNeg(
B);
4861 C = Builder.CreateFNeg(
C);
4868 if (Name[Name.size() - 5] ==
's')
4869 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4871 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4875 Rep = Builder.CreateFMA(
A,
B,
C);
4883 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4887 if (VecWidth == 128 && EltWidth == 32)
4888 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4889 else if (VecWidth == 256 && EltWidth == 32)
4890 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4891 else if (VecWidth == 128 && EltWidth == 64)
4892 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4893 else if (VecWidth == 256 && EltWidth == 64)
4894 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4900 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4901 Rep = Builder.CreateIntrinsic(IID,
Ops);
4902 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4903 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4904 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4905 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4906 bool IsMask3 = Name[11] ==
'3';
4907 bool IsMaskZ = Name[11] ==
'z';
4909 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4910 bool IsSubAdd = Name[3] ==
's';
4914 if (Name[Name.size() - 5] ==
's')
4915 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4917 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4922 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4924 Rep = Builder.CreateIntrinsic(IID,
Ops);
4933 Value *Odd = Builder.CreateCall(FMA,
Ops);
4934 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4935 Value *Even = Builder.CreateCall(FMA,
Ops);
4941 for (
int i = 0; i != NumElts; ++i)
4942 Idxs[i] = i + (i % 2) * NumElts;
4944 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4952 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4953 Name.starts_with(
"avx512.maskz.pternlog.")) {
4954 bool ZeroMask = Name[11] ==
'z';
4958 if (VecWidth == 128 && EltWidth == 32)
4959 IID = Intrinsic::x86_avx512_pternlog_d_128;
4960 else if (VecWidth == 256 && EltWidth == 32)
4961 IID = Intrinsic::x86_avx512_pternlog_d_256;
4962 else if (VecWidth == 512 && EltWidth == 32)
4963 IID = Intrinsic::x86_avx512_pternlog_d_512;
4964 else if (VecWidth == 128 && EltWidth == 64)
4965 IID = Intrinsic::x86_avx512_pternlog_q_128;
4966 else if (VecWidth == 256 && EltWidth == 64)
4967 IID = Intrinsic::x86_avx512_pternlog_q_256;
4968 else if (VecWidth == 512 && EltWidth == 64)
4969 IID = Intrinsic::x86_avx512_pternlog_q_512;
4975 Rep = Builder.CreateIntrinsic(IID, Args);
4979 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4980 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4981 bool ZeroMask = Name[11] ==
'z';
4982 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4985 if (VecWidth == 128 && !
High)
4986 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4987 else if (VecWidth == 256 && !
High)
4988 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4989 else if (VecWidth == 512 && !
High)
4990 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4991 else if (VecWidth == 128 &&
High)
4992 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4993 else if (VecWidth == 256 &&
High)
4994 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4995 else if (VecWidth == 512 &&
High)
4996 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5002 Rep = Builder.CreateIntrinsic(IID, Args);
5006 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5007 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5008 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5009 bool ZeroMask = Name[11] ==
'z';
5010 bool IndexForm = Name[17] ==
'i';
5012 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5013 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5014 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5015 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5016 bool ZeroMask = Name[11] ==
'z';
5017 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5020 if (VecWidth == 128 && !IsSaturating)
5021 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5022 else if (VecWidth == 256 && !IsSaturating)
5023 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5024 else if (VecWidth == 512 && !IsSaturating)
5025 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5026 else if (VecWidth == 128 && IsSaturating)
5027 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5028 else if (VecWidth == 256 && IsSaturating)
5029 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5030 else if (VecWidth == 512 && IsSaturating)
5031 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5041 if (Args[1]->
getType()->isVectorTy() &&
5044 ->isIntegerTy(32) &&
5045 Args[2]->
getType()->isVectorTy() &&
5048 ->isIntegerTy(32)) {
5049 Type *NewArgType =
nullptr;
5050 if (VecWidth == 128)
5052 else if (VecWidth == 256)
5054 else if (VecWidth == 512)
5060 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5061 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5064 Rep = Builder.CreateIntrinsic(IID, Args);
5068 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5069 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5070 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5071 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5072 bool ZeroMask = Name[11] ==
'z';
5073 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5076 if (VecWidth == 128 && !IsSaturating)
5077 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5078 else if (VecWidth == 256 && !IsSaturating)
5079 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5080 else if (VecWidth == 512 && !IsSaturating)
5081 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5082 else if (VecWidth == 128 && IsSaturating)
5083 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5084 else if (VecWidth == 256 && IsSaturating)
5085 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5086 else if (VecWidth == 512 && IsSaturating)
5087 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5097 if (Args[1]->
getType()->isVectorTy() &&
5100 ->isIntegerTy(32) &&
5101 Args[2]->
getType()->isVectorTy() &&
5104 ->isIntegerTy(32)) {
5105 Type *NewArgType =
nullptr;
5106 if (VecWidth == 128)
5108 else if (VecWidth == 256)
5110 else if (VecWidth == 512)
5116 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5117 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5120 Rep = Builder.CreateIntrinsic(IID, Args);
5124 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5125 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5126 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5128 if (Name[0] ==
'a' && Name.back() ==
'2')
5129 IID = Intrinsic::x86_addcarry_32;
5130 else if (Name[0] ==
'a' && Name.back() ==
'4')
5131 IID = Intrinsic::x86_addcarry_64;
5132 else if (Name[0] ==
's' && Name.back() ==
'2')
5133 IID = Intrinsic::x86_subborrow_32;
5134 else if (Name[0] ==
's' && Name.back() ==
'4')
5135 IID = Intrinsic::x86_subborrow_64;
5142 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5145 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5148 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5152 }
else if (Name.starts_with(
"avx512.mask.") &&
5155 }
else if (Name.starts_with(
"bmi.pdep.")) {
5157 }
else if (Name.starts_with(
"bmi.pext.")) {
5167 if (Name.starts_with(
"neon.bfcvt")) {
5168 if (Name.starts_with(
"neon.bfcvtn2")) {
5170 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5172 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5173 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5176 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5177 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5179 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5183 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5184 return Builder.CreateShuffleVector(
5187 return Builder.CreateFPTrunc(CI->
getOperand(0),
5190 }
else if (Name.starts_with(
"sve.fcvt")) {
5193 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5194 .
Case(
"sve.fcvtnt.bf16f32",
5195 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5207 if (Args[1]->
getType() != BadPredTy)
5210 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5211 BadPredTy, Args[1]);
5212 Args[1] = Builder.CreateIntrinsic(
5213 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5215 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5219 if (Name ==
"neon.vcvtfp2hf")
5220 return Builder.CreateBitCast(
5221 Builder.CreateFPTrunc(
5225 if (Name ==
"neon.vcvthf2fp")
5226 return Builder.CreateFPExt(
5227 Builder.CreateBitCast(
5237 if (Name ==
"mve.vctp64.old") {
5240 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5243 Value *C1 = Builder.CreateIntrinsic(
5244 Intrinsic::arm_mve_pred_v2i,
5246 return Builder.CreateIntrinsic(
5247 Intrinsic::arm_mve_pred_i2v,
5249 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5250 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5251 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5252 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5254 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5255 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5256 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5257 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5259 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5260 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5261 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5262 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5263 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5264 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5265 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5266 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5267 std::vector<Type *> Tys;
5271 case Intrinsic::arm_mve_mull_int_predicated:
5272 case Intrinsic::arm_mve_vqdmull_predicated:
5273 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5276 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5277 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5278 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5282 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5286 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5290 case Intrinsic::arm_cde_vcx1q_predicated:
5291 case Intrinsic::arm_cde_vcx1qa_predicated:
5292 case Intrinsic::arm_cde_vcx2q_predicated:
5293 case Intrinsic::arm_cde_vcx2qa_predicated:
5294 case Intrinsic::arm_cde_vcx3q_predicated:
5295 case Intrinsic::arm_cde_vcx3qa_predicated:
5302 std::vector<Value *>
Ops;
5304 Type *Ty =
Op->getType();
5305 if (Ty->getScalarSizeInBits() == 1) {
5306 Value *C1 = Builder.CreateIntrinsic(
5307 Intrinsic::arm_mve_pred_v2i,
5309 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5314 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5329 auto UpgradeLegacyWMMAIUIntrinsicCall =
5334 Args.push_back(Builder.getFalse());
5338 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5345 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5349 NewCall->copyMetadata(*CI);
5353 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5354 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5355 "intrinsic should have 7 arguments");
5358 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5360 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5361 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5362 "intrinsic should have 8 arguments");
5367 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5370 switch (
F->getIntrinsicID()) {
5373 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5374 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5375 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5376 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5377 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5378 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5393 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5396 F->getParent(),
F->getIntrinsicID(), Overloads);
5401 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5405 NewCall->copyMetadata(*CI);
5406 NewCall->takeName(CI);
5411 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5417 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5418 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5426 if (Name.starts_with(
"addrspacecast.nonnull")) {
5429 Value *ASC = Builder.CreateAddrSpaceCast(
5452 if (NumOperands < 3)
5465 bool IsVolatile =
false;
5469 if (NumOperands > 3)
5474 if (NumOperands > 5) {
5476 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5490 if (VT->getElementType()->isIntegerTy(16)) {
5493 Val = Builder.CreateBitCast(Val, AsBF16);
5501 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5503 unsigned AddrSpace = PtrTy->getAddressSpace();
5506 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5508 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5513 MDNode *RangeNotPrivate =
5516 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5522 return Builder.CreateBitCast(RMW, RetTy);
5543 return MAV->getMetadata();
5552 if (Name ==
"label") {
5554 }
else if (Name ==
"assign") {
5561 }
else if (Name ==
"declare") {
5565 }
else if (Name ==
"addr") {
5575 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5576 }
else if (Name ==
"value") {
5579 unsigned ExprOp = 2;
5594 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5602 int64_t OffsetVal =
Offset->getSExtValue();
5603 return Builder.CreateIntrinsic(OffsetVal >= 0
5604 ? Intrinsic::vector_splice_left
5605 : Intrinsic::vector_splice_right,
5607 {CI->getArgOperand(0), CI->getArgOperand(1),
5608 Builder.getInt32(std::abs(OffsetVal))});
5613 if (Name.starts_with(
"to.fp16")) {
5615 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5616 return Builder.CreateBitCast(Cast, CI->
getType());
5619 if (Name.starts_with(
"from.fp16")) {
5621 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5622 return Builder.CreateFPExt(Cast, CI->
getType());
5681 else if (Opcode == Instruction::ICmp)
5684 else if (Opcode == Instruction::FCmp)
5687 else if (Opcode == Instruction::Select)
5692 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5704 if (Defaults.empty())
5707 unsigned OldArgCount = CI->
arg_size();
5708 unsigned NewArgCount = NewFn->
arg_size();
5710 if (OldArgCount < FirstDefault)
5714 if (OldArgCount > NewArgCount)
5719 if (OldArgCount == NewArgCount) {
5731 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5732 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5733 "missing argument outside the default range");
5734 Type *ParamTy = NewFT->getParamType(Idx);
5739 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5745 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5776 if (!Name.consume_front(
"llvm."))
5779 bool IsX86 = Name.consume_front(
"x86.");
5780 bool IsNVVM = Name.consume_front(
"nvvm.");
5781 bool IsAArch64 = Name.consume_front(
"aarch64.");
5782 bool IsARM = Name.consume_front(
"arm.");
5783 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5784 bool IsDbg = Name.consume_front(
"dbg.");
5786 (Name.consume_front(
"experimental.vector.splice") ||
5787 Name.consume_front(
"vector.splice")) &&
5788 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5789 Value *Rep =
nullptr;
5791 if (!IsX86 && Name ==
"stackprotectorcheck") {
5793 }
else if (IsNVVM) {
5797 }
else if (IsAArch64) {
5801 }
else if (IsAMDGCN) {
5805 }
else if (IsOldSplice) {
5807 }
else if (Name.consume_front(
"convert.")) {
5809 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5824 const auto &DefaultCase = [&]() ->
void {
5832 "Unknown function for CallBase upgrade and isn't just a name change");
5840 "Return type must have changed");
5841 assert(OldST->getNumElements() ==
5843 "Must have same number of elements");
5846 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5849 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5850 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5851 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5872 case Intrinsic::arm_neon_vst1:
5873 case Intrinsic::arm_neon_vst2:
5874 case Intrinsic::arm_neon_vst3:
5875 case Intrinsic::arm_neon_vst4:
5876 case Intrinsic::arm_neon_vst2lane:
5877 case Intrinsic::arm_neon_vst3lane:
5878 case Intrinsic::arm_neon_vst4lane: {
5880 NewCall = Builder.CreateCall(NewFn, Args);
5883 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5884 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5885 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5890 NewCall = Builder.CreateCall(NewFn, Args);
5893 case Intrinsic::aarch64_sve_ld3_sret:
5894 case Intrinsic::aarch64_sve_ld4_sret:
5895 case Intrinsic::aarch64_sve_ld2_sret: {
5903 Name = Name.substr(5);
5910 unsigned MinElts = RetTy->getMinNumElements() /
N;
5912 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5914 for (
unsigned I = 0;
I <
N;
I++) {
5915 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5916 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5922 case Intrinsic::coro_end_async:
5923 case Intrinsic::coro_end: {
5925 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5927 NewCall = Builder.CreateCall(NewFn, Args);
5932 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5933 Value *InRamp = Builder.CreateCall(IsInRamp);
5943 case Intrinsic::vector_extract: {
5945 Name = Name.substr(5);
5946 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5951 unsigned MinElts = RetTy->getMinNumElements();
5954 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5958 case Intrinsic::vector_insert: {
5960 Name = Name.substr(5);
5961 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5965 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5970 NewCall = Builder.CreateCall(
5974 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5980 assert(
N > 1 &&
"Create is expected to be between 2-4");
5983 unsigned MinElts = RetTy->getMinNumElements() /
N;
5984 for (
unsigned I = 0;
I <
N;
I++) {
5986 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5993 case Intrinsic::arm_neon_bfdot:
5994 case Intrinsic::arm_neon_bfmmla:
5995 case Intrinsic::arm_neon_bfmlalb:
5996 case Intrinsic::arm_neon_bfmlalt:
5997 case Intrinsic::aarch64_neon_bfdot:
5998 case Intrinsic::aarch64_neon_bfmmla:
5999 case Intrinsic::aarch64_neon_bfmlalb:
6000 case Intrinsic::aarch64_neon_bfmlalt: {
6003 "Mismatch between function args and call args");
6004 size_t OperandWidth =
6006 assert((OperandWidth == 64 || OperandWidth == 128) &&
6007 "Unexpected operand width");
6009 auto Iter = CI->
args().begin();
6010 Args.push_back(*Iter++);
6011 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6012 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6013 NewCall = Builder.CreateCall(NewFn, Args);
6017 case Intrinsic::bitreverse:
6018 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6021 case Intrinsic::ctlz:
6022 case Intrinsic::cttz: {
6029 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6033 case Intrinsic::objectsize: {
6034 Value *NullIsUnknownSize =
6038 NewCall = Builder.CreateCall(
6043 case Intrinsic::ctpop:
6044 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6046 case Intrinsic::dbg_value: {
6048 Name = Name.substr(5);
6050 if (Name.starts_with(
"dbg.addr")) {
6064 if (
Offset->isNullValue()) {
6065 NewCall = Builder.CreateCall(
6074 case Intrinsic::ptr_annotation:
6082 NewCall = Builder.CreateCall(
6091 case Intrinsic::var_annotation:
6098 NewCall = Builder.CreateCall(
6107 case Intrinsic::riscv_aes32dsi:
6108 case Intrinsic::riscv_aes32dsmi:
6109 case Intrinsic::riscv_aes32esi:
6110 case Intrinsic::riscv_aes32esmi:
6111 case Intrinsic::riscv_sm4ks:
6112 case Intrinsic::riscv_sm4ed: {
6122 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6123 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6129 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6130 Value *Res = NewCall;
6132 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6138 case Intrinsic::nvvm_mapa_shared_cluster: {
6142 Value *Res = NewCall;
6143 Res = Builder.CreateAddrSpaceCast(
6150 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6152 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6154 Args[0] = Builder.CreateAddrSpaceCast(
6158 Args.push_back(Builder.getInt32(0));
6160 NewCall = Builder.CreateCall(NewFn, Args);
6166 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6171 for (
unsigned I = 0;
I < 4; ++
I)
6173 Args.push_back(Builder.getInt32(0));
6174 Args.push_back(Builder.getInt32(0));
6177 Args.push_back(Builder.getInt1(
false));
6178 Args.push_back(Builder.getInt32(0));
6180 NewCall = Builder.CreateCall(NewFn, Args);
6186 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6189 Args[0] = Builder.CreateAddrSpaceCast(
6192 NewCall = Builder.CreateCall(NewFn, Args);
6199#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6200 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6202#undef G2S_CLUSTER_CASE
6207 Args[0] = Builder.CreateAddrSpaceCast(
6213 Args.push_back(Builder.getInt32(0));
6215 NewCall = Builder.CreateCall(NewFn, Args);
6222#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6223 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6231 "expected only the trailing flag_valid_pattern to be missing");
6232 Args.push_back(Builder.getInt32(0));
6234 NewCall = Builder.CreateCall(NewFn, Args);
6240#undef NVVM_TMA_G2S_MODES
6243 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6244 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6245 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6246 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6247 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6248 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6249 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6250 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6252 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6256 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6257 NewCall = Builder.CreateCall(NewFn, Args);
6260 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6261 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6262 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6263 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6266 Builder.getFalse()});
6268 case Intrinsic::nvvm_mbarrier_init: {
6272 if (Args.size() == 2)
6273 Args.push_back(Builder.getInt32(0));
6274 NewCall = Builder.CreateCall(NewFn, Args);
6277 case Intrinsic::riscv_sha256sig0:
6278 case Intrinsic::riscv_sha256sig1:
6279 case Intrinsic::riscv_sha256sum0:
6280 case Intrinsic::riscv_sha256sum1:
6281 case Intrinsic::riscv_sm3p0:
6282 case Intrinsic::riscv_sm3p1: {
6289 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6291 NewCall = Builder.CreateCall(NewFn, Arg);
6293 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6300 case Intrinsic::x86_xop_vfrcz_ss:
6301 case Intrinsic::x86_xop_vfrcz_sd:
6302 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6305 case Intrinsic::x86_xop_vpermil2pd:
6306 case Intrinsic::x86_xop_vpermil2ps:
6307 case Intrinsic::x86_xop_vpermil2pd_256:
6308 case Intrinsic::x86_xop_vpermil2ps_256: {
6312 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6313 NewCall = Builder.CreateCall(NewFn, Args);
6317 case Intrinsic::x86_sse41_ptestc:
6318 case Intrinsic::x86_sse41_ptestz:
6319 case Intrinsic::x86_sse41_ptestnzc: {
6333 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6334 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6336 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6340 case Intrinsic::x86_rdtscp: {
6346 NewCall = Builder.CreateCall(NewFn);
6348 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6351 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6359 case Intrinsic::x86_sse41_insertps:
6360 case Intrinsic::x86_sse41_dppd:
6361 case Intrinsic::x86_sse41_dpps:
6362 case Intrinsic::x86_sse41_mpsadbw:
6363 case Intrinsic::x86_avx_dp_ps_256:
6364 case Intrinsic::x86_avx2_mpsadbw: {
6370 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6371 NewCall = Builder.CreateCall(NewFn, Args);
6375 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6376 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6377 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6378 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6379 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6380 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6386 NewCall = Builder.CreateCall(NewFn, Args);
6395 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6396 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6397 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6398 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6399 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6400 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6404 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6405 Args[1] = Builder.CreateBitCast(
6408 NewCall = Builder.CreateCall(NewFn, Args);
6409 Value *Res = Builder.CreateBitCast(
6417 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6418 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6419 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6423 Args[1] = Builder.CreateBitCast(
6425 Args[2] = Builder.CreateBitCast(
6428 NewCall = Builder.CreateCall(NewFn, Args);
6432 case Intrinsic::thread_pointer: {
6433 NewCall = Builder.CreateCall(NewFn, {});
6437 case Intrinsic::memcpy:
6438 case Intrinsic::memmove:
6439 case Intrinsic::memset: {
6455 NewCall = Builder.CreateCall(NewFn, Args);
6458 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6459 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6460 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6465 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6468 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6472 case Intrinsic::masked_load:
6473 case Intrinsic::masked_gather:
6474 case Intrinsic::masked_store:
6475 case Intrinsic::masked_scatter: {
6481 auto GetMaybeAlign = [](
Value *
Op) {
6483 uint64_t Val = CI->getZExtValue();
6491 auto GetAlign = [&](
Value *
Op) {
6500 case Intrinsic::masked_load:
6501 NewCall = Builder.CreateMaskedLoad(
6505 case Intrinsic::masked_gather:
6506 NewCall = Builder.CreateMaskedGather(
6512 case Intrinsic::masked_store:
6513 NewCall = Builder.CreateMaskedStore(
6517 case Intrinsic::masked_scatter:
6518 NewCall = Builder.CreateMaskedScatter(
6520 DL.getValueOrABITypeAlignment(
6534 case Intrinsic::lifetime_start:
6535 case Intrinsic::lifetime_end: {
6547 NewCall = Builder.CreateLifetimeStart(Ptr);
6549 NewCall = Builder.CreateLifetimeEnd(Ptr);
6558 case Intrinsic::x86_avx512_vpdpbusd_128:
6559 case Intrinsic::x86_avx512_vpdpbusd_256:
6560 case Intrinsic::x86_avx512_vpdpbusd_512:
6561 case Intrinsic::x86_avx512_vpdpbusds_128:
6562 case Intrinsic::x86_avx512_vpdpbusds_256:
6563 case Intrinsic::x86_avx512_vpdpbusds_512:
6564 case Intrinsic::x86_avx2_vpdpbssd_128:
6565 case Intrinsic::x86_avx2_vpdpbssd_256:
6566 case Intrinsic::x86_avx10_vpdpbssd_512:
6567 case Intrinsic::x86_avx2_vpdpbssds_128:
6568 case Intrinsic::x86_avx2_vpdpbssds_256:
6569 case Intrinsic::x86_avx10_vpdpbssds_512:
6570 case Intrinsic::x86_avx2_vpdpbsud_128:
6571 case Intrinsic::x86_avx2_vpdpbsud_256:
6572 case Intrinsic::x86_avx10_vpdpbsud_512:
6573 case Intrinsic::x86_avx2_vpdpbsuds_128:
6574 case Intrinsic::x86_avx2_vpdpbsuds_256:
6575 case Intrinsic::x86_avx10_vpdpbsuds_512:
6576 case Intrinsic::x86_avx2_vpdpbuud_128:
6577 case Intrinsic::x86_avx2_vpdpbuud_256:
6578 case Intrinsic::x86_avx10_vpdpbuud_512:
6579 case Intrinsic::x86_avx2_vpdpbuuds_128:
6580 case Intrinsic::x86_avx2_vpdpbuuds_256:
6581 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6586 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6587 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6589 NewCall = Builder.CreateCall(NewFn, Args);
6592 case Intrinsic::x86_avx512_vpdpwssd_128:
6593 case Intrinsic::x86_avx512_vpdpwssd_256:
6594 case Intrinsic::x86_avx512_vpdpwssd_512:
6595 case Intrinsic::x86_avx512_vpdpwssds_128:
6596 case Intrinsic::x86_avx512_vpdpwssds_256:
6597 case Intrinsic::x86_avx512_vpdpwssds_512:
6598 case Intrinsic::x86_avx2_vpdpwsud_128:
6599 case Intrinsic::x86_avx2_vpdpwsud_256:
6600 case Intrinsic::x86_avx10_vpdpwsud_512:
6601 case Intrinsic::x86_avx2_vpdpwsuds_128:
6602 case Intrinsic::x86_avx2_vpdpwsuds_256:
6603 case Intrinsic::x86_avx10_vpdpwsuds_512:
6604 case Intrinsic::x86_avx2_vpdpwusd_128:
6605 case Intrinsic::x86_avx2_vpdpwusd_256:
6606 case Intrinsic::x86_avx10_vpdpwusd_512:
6607 case Intrinsic::x86_avx2_vpdpwusds_128:
6608 case Intrinsic::x86_avx2_vpdpwusds_256:
6609 case Intrinsic::x86_avx10_vpdpwusds_512:
6610 case Intrinsic::x86_avx2_vpdpwuud_128:
6611 case Intrinsic::x86_avx2_vpdpwuud_256:
6612 case Intrinsic::x86_avx10_vpdpwuud_512:
6613 case Intrinsic::x86_avx2_vpdpwuuds_128:
6614 case Intrinsic::x86_avx2_vpdpwuuds_256:
6615 case Intrinsic::x86_avx10_vpdpwuuds_512:
6620 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6621 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6623 NewCall = Builder.CreateCall(NewFn, Args);
6626 assert(NewCall &&
"Should have either set this variable or returned through "
6627 "the default case");
6634 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6648 F->eraseFromParent();
6654 if (NumOperands == 0)
6662 if (NumOperands == 3) {
6666 Metadata *Elts2[] = {ScalarType, ScalarType,
6682 if (NumOperands == 0 || NumOperands % 3 != 0)
6687 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6692 if (Upgraded ==
Tag)
6702 if (
Opc != Instruction::BitCast)
6706 Type *SrcTy = V->getType();
6723 if (
Opc != Instruction::BitCast)
6726 Type *SrcTy =
C->getType();
6743 if (Flag.getNumOperands() < 3)
6744 return std::nullopt;
6746 return Name->getString();
6747 return std::nullopt;
6761 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6762 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6763 if (auto Name = getModuleFlagNameSafely(*Flag))
6764 return *Name ==
"Debug Info Version";
6767 if (OpIt != ModFlags->op_end()) {
6768 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6775 bool BrokenDebugInfo =
false;
6778 if (!BrokenDebugInfo)
6784 M.getContext().diagnose(Diag);
6791 M.getContext().diagnose(DiagVersion);
6801 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6804 if (
F->hasFnAttribute(Attr)) {
6807 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6809 auto [Part, Rest] = S.
split(
',');
6815 const unsigned Dim = DimC -
'x';
6816 assert(Dim < 3 &&
"Unexpected dim char");
6826 F->addFnAttr(Attr, NewAttr);
6830 return S ==
"x" || S ==
"y" || S ==
"z";
6835 if (
K ==
"kernel") {
6847 const unsigned Idx = (AlignIdxValuePair >> 16);
6848 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6853 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6858 if (
K ==
"minctasm") {
6863 if (
K ==
"maxnreg") {
6868 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6872 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6876 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6880 if (
K ==
"grid_constant") {
6895 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6902 if (!SeenNodes.
insert(MD).second)
6909 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6916 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6918 const MDOperand &V = MD->getOperand(j + 1);
6924 if (NewOperands.
size() > 1)
6937 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6938 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6939 if (ModRetainReleaseMarker) {
6945 ID->getString().split(ValueComp,
"#");
6946 if (ValueComp.
size() == 2) {
6947 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6951 M.eraseNamedMetadata(ModRetainReleaseMarker);
6962 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6988 bool InvalidCast =
false;
6990 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7003 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7005 Args.push_back(Arg);
7012 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7017 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7030 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7038 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7039 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7040 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7041 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7042 {
"objc_autoreleaseReturnValue",
7043 llvm::Intrinsic::objc_autoreleaseReturnValue},
7044 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7045 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7046 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7047 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7048 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7049 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7050 {
"objc_release", llvm::Intrinsic::objc_release},
7051 {
"objc_retain", llvm::Intrinsic::objc_retain},
7052 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7053 {
"objc_retainAutoreleaseReturnValue",
7054 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7055 {
"objc_retainAutoreleasedReturnValue",
7056 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7057 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7058 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7059 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7060 {
"objc_unsafeClaimAutoreleasedReturnValue",
7061 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7062 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7063 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7064 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7065 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7066 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7067 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7068 {
"objc_arc_annotation_topdown_bbstart",
7069 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7070 {
"objc_arc_annotation_topdown_bbend",
7071 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7072 {
"objc_arc_annotation_bottomup_bbstart",
7073 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7074 {
"objc_arc_annotation_bottomup_bbend",
7075 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7077 for (
auto &
I : RuntimeFuncs)
7078 UpgradeToIntrinsic(
I.first,
I.second);
7102 std::optional<bool> UseAddressDisc;
7105 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7106 for (
const MDNode *Flag : ModFlags->operands()) {
7108 if (Name && (*Name ==
"ptrauth-init-fini" ||
7109 *Name ==
"ptrauth-init-fini-address-discrimination"))
7114 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7115 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7116 constexpr unsigned ExpectedAddressMarker = 1;
7119 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7122 bool HasAddressDisc;
7123 if (!CPA->hasAddressDiscriminator())
7124 HasAddressDisc =
false;
7125 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7126 HasAddressDisc =
true;
7130 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7133 UseAddressDisc = HasAddressDisc;
7134 return CPA->getPointer();
7138 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7141 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7143 if (!GV || !GV->hasInitializer())
7147 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7150 std::vector<Constant *> NewStructors;
7151 NewStructors.reserve(OldStructorsArray->getNumOperands());
7153 for (
Use &U : OldStructorsArray->operands()) {
7162 Func = UpgradeSinglePointer(Func);
7166 NewStructors.push_back(
7175 if (GlobalArraysToUpgrade.
empty())
7177 assert(UseAddressDisc.has_value());
7179 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7180 GV->setInitializer(NewInit);
7183 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7193 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7197 bool HasObjCFlag =
false, HasClassProperties =
false;
7198 bool HasSwiftVersionFlag =
false;
7199 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7206 if (
Op->getNumOperands() != 3)
7220 if (ID->getString() ==
"Objective-C Image Info Version")
7222 if (ID->getString() ==
"Objective-C Class Properties")
7223 HasClassProperties =
true;
7225 if (ID->getString() ==
"PIC Level") {
7226 if (
auto *Behavior =
7228 uint64_t V = Behavior->getLimitedValue();
7234 if (ID->getString() ==
"PIE Level")
7235 if (
auto *Behavior =
7242 if (ID->getString() ==
"branch-target-enforcement" ||
7243 ID->getString().starts_with(
"sign-return-address")) {
7244 if (
auto *Behavior =
7250 Op->getOperand(1),
Op->getOperand(2)};
7260 if (ID->getString() ==
"Objective-C Image Info Section") {
7263 Value->getString().split(ValueComp,
" ");
7264 if (ValueComp.
size() != 1) {
7265 std::string NewValue;
7266 for (
auto &S : ValueComp)
7267 NewValue += S.str();
7278 if (ID->getString() ==
"Objective-C Garbage Collection") {
7281 assert(Md->getValue() &&
"Expected non-empty metadata");
7282 auto Type = Md->getValue()->getType();
7285 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7286 if ((Val & 0xff) != Val) {
7287 HasSwiftVersionFlag =
true;
7288 SwiftABIVersion = (Val & 0xff00) >> 8;
7289 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7290 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7301 if (ID->getString() ==
"amdgpu_code_object_version") {
7304 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7313 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7342 if (HasObjCFlag && !HasClassProperties) {
7348 if (HasSwiftVersionFlag) {
7352 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7354 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7362 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7366 auto MatchesVersion = [](
const MDNode *
Op) {
7367 return Op->getNumOperands() >= 3 &&
7381 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7382 assert(
Op->getNumOperands() >= 2 &&
7383 "Expected at least 2 operands - name and linkage type");
7395 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7406 auto TrimSpaces = [](
StringRef Section) -> std::string {
7408 Section.split(Components,
',');
7413 for (
auto Component : Components)
7414 OS <<
',' << Component.trim();
7419 for (
auto &GV : M.globals()) {
7420 if (!GV.hasSection())
7425 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7430 GV.setSection(TrimSpaces(Section));
7446struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7447 StrictFPUpgradeVisitor() =
default;
7450 if (!
Call.isStrictFP())
7456 Call.removeFnAttr(Attribute::StrictFP);
7457 Call.addFnAttr(Attribute::NoBuiltin);
7462struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7463 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7464 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7466 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7481 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7482 StrictFPUpgradeVisitor SFPV;
7487 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7488 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7489 for (
auto &Arg :
F.args())
7491 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7493 bool AddingAttrs =
false, RemovingAttrs =
false;
7494 AttrBuilder AttrsToAdd(
F.getContext());
7499 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7500 A.isValid() &&
A.isStringAttribute()) {
7501 F.setSection(
A.getValueAsString());
7503 RemovingAttrs =
true;
7507 A.isValid() &&
A.isStringAttribute()) {
7510 AddingAttrs = RemovingAttrs =
true;
7513 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7514 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7516 RemovingAttrs =
true;
7517 if (
A.getValueAsString() ==
"true") {
7518 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7527 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7530 if (
A.getValueAsBool()) {
7531 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7537 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7538 RemovingAttrs =
true;
7545 bool HandleDenormalMode =
false;
7547 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7550 DenormalFPMath = ParsedMode;
7552 AddingAttrs = RemovingAttrs =
true;
7553 HandleDenormalMode =
true;
7557 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7561 DenormalFPMathF32 = ParsedMode;
7563 AddingAttrs = RemovingAttrs =
true;
7564 HandleDenormalMode =
true;
7568 if (HandleDenormalMode)
7569 AttrsToAdd.addDenormalFPEnvAttr(
7573 F.removeFnAttrs(AttrsToRemove);
7576 F.addFnAttrs(AttrsToAdd);
7582 if (!
F.hasFnAttribute(FnAttrName))
7583 F.addFnAttr(FnAttrName,
Value);
7590 if (!
F.hasFnAttribute(FnAttrName)) {
7592 F.addFnAttr(FnAttrName);
7594 auto A =
F.getFnAttribute(FnAttrName);
7595 if (
"false" ==
A.getValueAsString())
7596 F.removeFnAttr(FnAttrName);
7597 else if (
"true" ==
A.getValueAsString()) {
7598 F.removeFnAttr(FnAttrName);
7599 F.addFnAttr(FnAttrName);
7605 Triple T(M.getTargetTriple());
7606 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7609 uint64_t BTEValue = 0;
7610 uint64_t BPPLRValue = 0;
7611 uint64_t GCSValue = 0;
7612 uint64_t SRAValue = 0;
7613 uint64_t SRAALLValue = 0;
7614 uint64_t SRABKeyValue = 0;
7616 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7620 if (
Op->getNumOperands() != 3)
7629 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7630 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7631 : IDStr ==
"guarded-control-stack" ? &GCSValue
7632 : IDStr ==
"sign-return-address" ? &SRAValue
7633 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7634 : IDStr ==
"sign-return-address-with-bkey"
7640 *ValPtr = CI->getZExtValue();
7646 bool BTE = BTEValue == 1;
7647 bool BPPLR = BPPLRValue == 1;
7648 bool GCS = GCSValue == 1;
7649 bool SRA = SRAValue == 1;
7652 if (SRA && SRAALLValue == 1)
7653 SignTypeValue =
"all";
7656 if (SRA && SRABKeyValue == 1)
7657 SignKeyValue =
"b_key";
7659 for (
Function &
F : M.getFunctionList()) {
7660 if (
F.isDeclaration())
7667 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7668 A.isValid() &&
"none" ==
A.getValueAsString()) {
7669 F.removeFnAttr(
"sign-return-address");
7670 F.removeFnAttr(
"sign-return-address-key");
7686 if (SRAALLValue == 1)
7688 if (SRABKeyValue == 1)
7715 if (
T->getNumOperands() < 1)
7720 if (S->getString().starts_with(
"llvm.vectorizer."))
7726 StringRef OldPrefix =
"llvm.vectorizer.";
7729 if (OldTag ==
"llvm.vectorizer.unroll")
7741 if (
T->getNumOperands() < 1)
7753 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7766 Ops.reserve(
T->getNumOperands());
7767 Ops.push_back(NewTag);
7768 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7769 Ops.push_back(
T->getOperand(
I));
7786 if (
T->isDistinct()) {
7787 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7799 Ops.reserve(
T->getNumOperands());
7810 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7811 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7812 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7815 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7817 auto I =
DL.find(
"-n64-");
7819 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7824 std::string Res =
DL.str();
7827 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7828 Res.append(Res.empty() ?
"G1" :
"-G1");
7836 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7837 Res.append(
"-ni:7:8:9");
7839 if (
DL.ends_with(
"ni:7"))
7841 if (
DL.ends_with(
"ni:7:8"))
7846 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7847 Res.append(
"-p7:160:256:256:32");
7848 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7849 Res.append(
"-p8:128:128:128:48");
7850 constexpr StringRef OldP8(
"-p8:128:128-");
7851 if (
DL.contains(OldP8))
7852 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7853 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7854 Res.append(
"-p9:192:256:256:32");
7859 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7860 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7861 Res.append((
"-" + AS +
":32:32").str());
7866 if (!
DL.contains(
"m:e"))
7867 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7872 if (
T.isSystemZ() && !
DL.empty()) {
7874 if (!
DL.contains(
"-S64"))
7875 return "E-S64" +
DL.drop_front(1).str();
7879 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7882 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7883 if (!
DL.contains(AddrSpaces)) {
7885 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7886 if (R.match(Res, &
Groups))
7892 if (
T.isAArch64()) {
7894 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7895 Res.append(
"-Fn32");
7896 AddPtr32Ptr64AddrSpaces();
7900 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7904 std::string I64 =
"-i64:64";
7905 std::string I128 =
"-i128:128";
7907 size_t Pos = Res.find(I64);
7908 if (Pos !=
size_t(-1))
7909 Res.insert(Pos + I64.size(), I128);
7913 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7914 size_t Pos = Res.find(
"-S128");
7917 Res.insert(Pos,
"-f64:32:64");
7922 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
7924 size_t Pos = Res.
find(p3232);
7926 Res.insert(Pos + p3232.
size(),
"-Fi8");
7932 AddPtr32Ptr64AddrSpaces();
7940 if (!
T.isOSIAMCU()) {
7941 std::string I128 =
"-i128:128";
7944 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7945 if (R.match(Res, &
Groups))
7953 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7955 auto I =
Ref.find(
"-f80:32-");
7957 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7965 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7968 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7969 B.removeAttribute(
"no-frame-pointer-elim");
7971 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7973 if (FramePointer !=
"all")
7974 FramePointer =
"non-leaf";
7975 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7977 if (!FramePointer.
empty())
7978 B.addAttribute(
"frame-pointer", FramePointer);
7980 A =
B.getAttribute(
"null-pointer-is-valid");
7983 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7984 B.removeAttribute(
"null-pointer-is-valid");
7985 if (NullPointerIsValid)
7986 B.addAttribute(Attribute::NullPointerIsValid);
7989 A =
B.getAttribute(
"uniform-work-group-size");
7993 bool IsTrue = Val ==
"true";
7994 B.removeAttribute(
"uniform-work-group-size");
7996 B.addAttribute(
"uniform-work-group-size");
8007 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
static unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID, SmallVectorImpl< Type * > &OverloadTys)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFAddUpgrade(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 upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static 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 Intrinsic::ID shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
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 Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F, StringRef Name)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_CASE(ID_SUFFIX, NAME)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static FCmpInst::Predicate getVPFPPredicateFromMD(const Value *Op)
static GCRegistry::Add< ShadowStackGC > C("shadow-stack", "Very portable GC for uncooperative code generators")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
This file contains constants used for implementing Dwarf debug support.
Module.h This file contains the declarations for the Module class.
const AbstractManglingParser< Derived, Alloc >::OperatorInfo AbstractManglingParser< Derived, Alloc >::Ops[]
static bool isZero(Value *V, const DataLayout &DL, DominatorTree *DT, AssumptionCache *AC)
NVPTX address space definition.
This file contains the definitions of the enumerations and flags associated with NVVM Intrinsics,...
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
static SymbolRef::Type getType(const Symbol *Sym)
LocallyHashedType DenseMapInfo< LocallyHashedType >::Empty
static const X86InstrFMA3Group Groups[]
Class for arbitrary precision integers.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
Class to represent array types.
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Type * getElementType() const
an instruction that atomically reads a memory location, combines it with another value,...
void setVolatile(bool V)
Specify whether this is a volatile RMW or not.
BinOp
This enumeration lists the possible modifications atomicrmw can make.
@ USubCond
Subtract only if no unsigned overflow.
@ Min
*p = old <signed v ? old : v
@ USubSat
*p = usub.sat(old, v) usub.sat matches the behavior of llvm.usub.sat.
@ UIncWrap
Increment one up to a maximum value.
@ Max
*p = old >signed v ? old : v
@ FMin
*p = minnum(old, v) minnum matches the behavior of llvm.minnum.
@ FMax
*p = maxnum(old, v) maxnum matches the behavior of llvm.maxnum.
@ UDecWrap
Decrement one until a minimum value or zero.
bool isFloatingPointOperation() const
This class stores enough information to efficiently remove some attributes from an existing AttrBuild...
AttributeMask & addAttribute(Attribute::AttrKind Val)
Add an attribute to the mask.
Functions, function parameters, and return types can have attributes to indicate how they should be t...
static LLVM_ABI Attribute getWithStackAlignment(LLVMContext &Context, Align Alignment)
static LLVM_ABI Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
void setCallingConv(CallingConv::ID CC)
LLVM_ABI void getOperandBundlesAsDefs(SmallVectorImpl< OperandBundleDef > &Defs) const
Return the list of operand bundles attached to this instruction as a vector of OperandBundleDefs.
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
CallingConv::ID getCallingConv() const
Value * getCalledOperand() const
void setAttributes(AttributeList A)
Set the attributes for this call.
Value * getArgOperand(unsigned i) const
FunctionType * getFunctionType() const
LLVM_ABI Intrinsic::ID getIntrinsicID() const
Returns the intrinsic ID of the intrinsic called or Intrinsic::not_intrinsic if the called function i...
iterator_range< User::op_iterator > args()
Iteration adapter for range-for loops.
void setCalledOperand(Value *V)
unsigned arg_size() const
AttributeList getAttributes() const
Return the attributes for this call.
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
This class represents a function call, abstracting a target machine's calling convention.
void setTailCallKind(TailCallKind TCK)
static LLVM_ABI CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", InsertPosition InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
static LLVM_ABI bool castIsValid(Instruction::CastOps op, Type *SrcTy, Type *DstTy)
This method can be used to determine if a cast from SrcTy to DstTy using Opcode op is valid or not.
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
unsigned getNumParams() const
Return the number of fixed parameters this function type requires.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
op_iterator op_end() const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
op_iterator op_begin() const
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
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 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 FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
LLVM_ABI MDNode * UpgradeTBAAStructNode(MDNode &TBAAStructNode)
If the given !tbaa.struct node has old-style scalar field tags, return an equivalent node with each f...
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
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...
bool is_contained(R &&Range, const E &Element)
Returns true if Element is found in Range.
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.