35#include "llvm/IR/IntrinsicsAArch64.h"
36#include "llvm/IR/IntrinsicsAMDGPU.h"
37#include "llvm/IR/IntrinsicsARM.h"
38#include "llvm/IR/IntrinsicsNVPTX.h"
39#include "llvm/IR/IntrinsicsRISCV.h"
40#include "llvm/IR/IntrinsicsWebAssembly.h"
41#include "llvm/IR/IntrinsicsX86.h"
65#define DEBUG_TYPE "auto-upgrade"
69 cl::desc(
"Disable autoupgrade of debug info"));
88 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
103 Type *LastArgType =
F->getFunctionType()->getParamType(
104 F->getFunctionType()->getNumParams() - 1);
119 if (
F->getReturnType()->isVectorTy())
132 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
133 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
150 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
151 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
165 if (
F->getReturnType()->getScalarType()->isBFloatTy())
175 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
189 if (Name.consume_front(
"avx."))
190 return (Name.starts_with(
"blend.p") ||
191 Name ==
"cvt.ps2.pd.256" ||
192 Name ==
"cvtdq2.pd.256" ||
193 Name ==
"cvtdq2.ps.256" ||
194 Name.starts_with(
"movnt.") ||
195 Name.starts_with(
"sqrt.p") ||
196 Name.starts_with(
"storeu.") ||
197 Name.starts_with(
"vbroadcast.s") ||
198 Name.starts_with(
"vbroadcastf128") ||
199 Name.starts_with(
"vextractf128.") ||
200 Name.starts_with(
"vinsertf128.") ||
201 Name.starts_with(
"vperm2f128.") ||
202 Name.starts_with(
"vpermil."));
204 if (Name.consume_front(
"avx2."))
205 return (Name ==
"movntdqa" ||
206 Name.starts_with(
"pabs.") ||
207 Name.starts_with(
"padds.") ||
208 Name.starts_with(
"paddus.") ||
209 Name.starts_with(
"pblendd.") ||
211 Name.starts_with(
"pbroadcast") ||
212 Name.starts_with(
"pcmpeq.") ||
213 Name.starts_with(
"pcmpgt.") ||
214 Name.starts_with(
"pmax") ||
215 Name.starts_with(
"pmin") ||
216 Name.starts_with(
"pmovsx") ||
217 Name.starts_with(
"pmovzx") ||
218 Name.starts_with(
"pmulh.w") ||
219 Name.starts_with(
"pmulhu.w") ||
221 Name ==
"pmulu.dq" ||
222 Name.starts_with(
"psll.dq") ||
223 Name.starts_with(
"psrl.dq") ||
224 Name.starts_with(
"psubs.") ||
225 Name.starts_with(
"psubus.") ||
226 Name.starts_with(
"vbroadcast") ||
227 Name ==
"vbroadcasti128" ||
228 Name ==
"vextracti128" ||
229 Name ==
"vinserti128" ||
230 Name ==
"vperm2i128");
232 if (Name.consume_front(
"avx512.")) {
233 if (Name.consume_front(
"mask."))
235 return (Name.starts_with(
"add.p") ||
236 Name.starts_with(
"and.") ||
237 Name.starts_with(
"andn.") ||
238 Name.starts_with(
"broadcast.s") ||
239 Name.starts_with(
"broadcastf32x4.") ||
240 Name.starts_with(
"broadcastf32x8.") ||
241 Name.starts_with(
"broadcastf64x2.") ||
242 Name.starts_with(
"broadcastf64x4.") ||
243 Name.starts_with(
"broadcasti32x4.") ||
244 Name.starts_with(
"broadcasti32x8.") ||
245 Name.starts_with(
"broadcasti64x2.") ||
246 Name.starts_with(
"broadcasti64x4.") ||
247 Name.starts_with(
"cmp.b") ||
248 Name.starts_with(
"cmp.d") ||
249 Name.starts_with(
"cmp.q") ||
250 Name.starts_with(
"cmp.w") ||
251 Name.starts_with(
"compress.b") ||
252 Name.starts_with(
"compress.d") ||
253 Name.starts_with(
"compress.p") ||
254 Name.starts_with(
"compress.q") ||
255 Name.starts_with(
"compress.store.") ||
256 Name.starts_with(
"compress.w") ||
257 Name.starts_with(
"conflict.") ||
258 Name.starts_with(
"cvtdq2pd.") ||
259 Name.starts_with(
"cvtdq2ps.") ||
260 Name ==
"cvtpd2dq.256" ||
261 Name ==
"cvtpd2ps.256" ||
262 Name ==
"cvtps2pd.128" ||
263 Name ==
"cvtps2pd.256" ||
264 Name.starts_with(
"cvtqq2pd.") ||
265 Name ==
"cvtqq2ps.256" ||
266 Name ==
"cvtqq2ps.512" ||
267 Name ==
"cvttpd2dq.256" ||
268 Name ==
"cvttps2dq.128" ||
269 Name ==
"cvttps2dq.256" ||
270 Name.starts_with(
"cvtudq2pd.") ||
271 Name.starts_with(
"cvtudq2ps.") ||
272 Name.starts_with(
"cvtuqq2pd.") ||
273 Name ==
"cvtuqq2ps.256" ||
274 Name ==
"cvtuqq2ps.512" ||
275 Name.starts_with(
"dbpsadbw.") ||
276 Name.starts_with(
"div.p") ||
277 Name.starts_with(
"expand.b") ||
278 Name.starts_with(
"expand.d") ||
279 Name.starts_with(
"expand.load.") ||
280 Name.starts_with(
"expand.p") ||
281 Name.starts_with(
"expand.q") ||
282 Name.starts_with(
"expand.w") ||
283 Name.starts_with(
"fpclass.p") ||
284 Name.starts_with(
"insert") ||
285 Name.starts_with(
"load.") ||
286 Name.starts_with(
"loadu.") ||
287 Name.starts_with(
"lzcnt.") ||
288 Name.starts_with(
"max.p") ||
289 Name.starts_with(
"min.p") ||
290 Name.starts_with(
"movddup") ||
291 Name.starts_with(
"move.s") ||
292 Name.starts_with(
"movshdup") ||
293 Name.starts_with(
"movsldup") ||
294 Name.starts_with(
"mul.p") ||
295 Name.starts_with(
"or.") ||
296 Name.starts_with(
"pabs.") ||
297 Name.starts_with(
"packssdw.") ||
298 Name.starts_with(
"packsswb.") ||
299 Name.starts_with(
"packusdw.") ||
300 Name.starts_with(
"packuswb.") ||
301 Name.starts_with(
"padd.") ||
302 Name.starts_with(
"padds.") ||
303 Name.starts_with(
"paddus.") ||
304 Name.starts_with(
"palignr.") ||
305 Name.starts_with(
"pand.") ||
306 Name.starts_with(
"pandn.") ||
307 Name.starts_with(
"pavg") ||
308 Name.starts_with(
"pbroadcast") ||
309 Name.starts_with(
"pcmpeq.") ||
310 Name.starts_with(
"pcmpgt.") ||
311 Name.starts_with(
"perm.df.") ||
312 Name.starts_with(
"perm.di.") ||
313 Name.starts_with(
"permvar.") ||
314 Name.starts_with(
"pmaddubs.w.") ||
315 Name.starts_with(
"pmaddw.d.") ||
316 Name.starts_with(
"pmax") ||
317 Name.starts_with(
"pmin") ||
318 Name ==
"pmov.qd.256" ||
319 Name ==
"pmov.qd.512" ||
320 Name ==
"pmov.wb.256" ||
321 Name ==
"pmov.wb.512" ||
322 Name.starts_with(
"pmovsx") ||
323 Name.starts_with(
"pmovzx") ||
324 Name.starts_with(
"pmul.dq.") ||
325 Name.starts_with(
"pmul.hr.sw.") ||
326 Name.starts_with(
"pmulh.w.") ||
327 Name.starts_with(
"pmulhu.w.") ||
328 Name.starts_with(
"pmull.") ||
329 Name.starts_with(
"pmultishift.qb.") ||
330 Name.starts_with(
"pmulu.dq.") ||
331 Name.starts_with(
"por.") ||
332 Name.starts_with(
"prol.") ||
333 Name.starts_with(
"prolv.") ||
334 Name.starts_with(
"pror.") ||
335 Name.starts_with(
"prorv.") ||
336 Name.starts_with(
"pshuf.b.") ||
337 Name.starts_with(
"pshuf.d.") ||
338 Name.starts_with(
"pshufh.w.") ||
339 Name.starts_with(
"pshufl.w.") ||
340 Name.starts_with(
"psll.d") ||
341 Name.starts_with(
"psll.q") ||
342 Name.starts_with(
"psll.w") ||
343 Name.starts_with(
"pslli") ||
344 Name.starts_with(
"psllv") ||
345 Name.starts_with(
"psra.d") ||
346 Name.starts_with(
"psra.q") ||
347 Name.starts_with(
"psra.w") ||
348 Name.starts_with(
"psrai") ||
349 Name.starts_with(
"psrav") ||
350 Name.starts_with(
"psrl.d") ||
351 Name.starts_with(
"psrl.q") ||
352 Name.starts_with(
"psrl.w") ||
353 Name.starts_with(
"psrli") ||
354 Name.starts_with(
"psrlv") ||
355 Name.starts_with(
"psub.") ||
356 Name.starts_with(
"psubs.") ||
357 Name.starts_with(
"psubus.") ||
358 Name.starts_with(
"pternlog.") ||
359 Name.starts_with(
"punpckh") ||
360 Name.starts_with(
"punpckl") ||
361 Name.starts_with(
"pxor.") ||
362 Name.starts_with(
"shuf.f") ||
363 Name.starts_with(
"shuf.i") ||
364 Name.starts_with(
"shuf.p") ||
365 Name.starts_with(
"sqrt.p") ||
366 Name.starts_with(
"store.b.") ||
367 Name.starts_with(
"store.d.") ||
368 Name.starts_with(
"store.p") ||
369 Name.starts_with(
"store.q.") ||
370 Name.starts_with(
"store.w.") ||
371 Name ==
"store.ss" ||
372 Name.starts_with(
"storeu.") ||
373 Name.starts_with(
"sub.p") ||
374 Name.starts_with(
"ucmp.") ||
375 Name.starts_with(
"unpckh.") ||
376 Name.starts_with(
"unpckl.") ||
377 Name.starts_with(
"valign.") ||
378 Name ==
"vcvtph2ps.128" ||
379 Name ==
"vcvtph2ps.256" ||
380 Name.starts_with(
"vextract") ||
381 Name.starts_with(
"vfmadd.") ||
382 Name.starts_with(
"vfmaddsub.") ||
383 Name.starts_with(
"vfnmadd.") ||
384 Name.starts_with(
"vfnmsub.") ||
385 Name.starts_with(
"vpdpbusd.") ||
386 Name.starts_with(
"vpdpbusds.") ||
387 Name.starts_with(
"vpdpwssd.") ||
388 Name.starts_with(
"vpdpwssds.") ||
389 Name.starts_with(
"vpermi2var.") ||
390 Name.starts_with(
"vpermil.p") ||
391 Name.starts_with(
"vpermilvar.") ||
392 Name.starts_with(
"vpermt2var.") ||
393 Name.starts_with(
"vpmadd52") ||
394 Name.starts_with(
"vpshld.") ||
395 Name.starts_with(
"vpshldv.") ||
396 Name.starts_with(
"vpshrd.") ||
397 Name.starts_with(
"vpshrdv.") ||
398 Name.starts_with(
"vpshufbitqmb.") ||
399 Name.starts_with(
"xor."));
401 if (Name.consume_front(
"mask3."))
403 return (Name.starts_with(
"vfmadd.") ||
404 Name.starts_with(
"vfmaddsub.") ||
405 Name.starts_with(
"vfmsub.") ||
406 Name.starts_with(
"vfmsubadd.") ||
407 Name.starts_with(
"vfnmsub."));
409 if (Name.consume_front(
"maskz."))
411 return (Name.starts_with(
"pternlog.") ||
412 Name.starts_with(
"vfmadd.") ||
413 Name.starts_with(
"vfmaddsub.") ||
414 Name.starts_with(
"vpdpbusd.") ||
415 Name.starts_with(
"vpdpbusds.") ||
416 Name.starts_with(
"vpdpwssd.") ||
417 Name.starts_with(
"vpdpwssds.") ||
418 Name.starts_with(
"vpermt2var.") ||
419 Name.starts_with(
"vpmadd52") ||
420 Name.starts_with(
"vpshldv.") ||
421 Name.starts_with(
"vpshrdv."));
424 return (Name ==
"movntdqa" ||
425 Name ==
"pmul.dq.512" ||
426 Name ==
"pmulu.dq.512" ||
427 Name.starts_with(
"broadcastm") ||
428 Name.starts_with(
"cmp.p") ||
429 Name.starts_with(
"cvtb2mask.") ||
430 Name.starts_with(
"cvtd2mask.") ||
431 Name.starts_with(
"cvtmask2") ||
432 Name.starts_with(
"cvtq2mask.") ||
433 Name ==
"cvtusi2sd" ||
434 Name.starts_with(
"cvtw2mask.") ||
439 Name ==
"kortestc.w" ||
440 Name ==
"kortestz.w" ||
441 Name.starts_with(
"kunpck") ||
444 Name.starts_with(
"padds.") ||
445 Name.starts_with(
"pbroadcast") ||
446 Name.starts_with(
"pmulh.w") ||
447 Name.starts_with(
"pmulhu.w") ||
448 Name.starts_with(
"prol") ||
449 Name.starts_with(
"pror") ||
450 Name.starts_with(
"psll.dq") ||
451 Name.starts_with(
"psrl.dq") ||
452 Name.starts_with(
"psubs.") ||
453 Name.starts_with(
"ptestm") ||
454 Name.starts_with(
"ptestnm") ||
455 Name.starts_with(
"storent.") ||
456 Name.starts_with(
"vbroadcast.s") ||
457 Name.starts_with(
"vpshld.") ||
458 Name.starts_with(
"vpshrd."));
461 if (Name.consume_front(
"fma."))
462 return (Name.starts_with(
"vfmadd.") ||
463 Name.starts_with(
"vfmsub.") ||
464 Name.starts_with(
"vfmsubadd.") ||
465 Name.starts_with(
"vfnmadd.") ||
466 Name.starts_with(
"vfnmsub."));
468 if (Name.consume_front(
"fma4."))
469 return Name.starts_with(
"vfmadd.s");
471 if (Name.consume_front(
"sse."))
472 return (Name ==
"add.ss" ||
473 Name ==
"cvtsi2ss" ||
474 Name ==
"cvtsi642ss" ||
477 Name.starts_with(
"sqrt.p") ||
479 Name.starts_with(
"storeu.") ||
482 if (Name.consume_front(
"sse2."))
483 return (Name ==
"add.sd" ||
484 Name ==
"cvtdq2pd" ||
485 Name ==
"cvtdq2ps" ||
486 Name ==
"cvtps2pd" ||
487 Name ==
"cvtsi2sd" ||
488 Name ==
"cvtsi642sd" ||
489 Name ==
"cvtss2sd" ||
492 Name.starts_with(
"padds.") ||
493 Name.starts_with(
"paddus.") ||
494 Name.starts_with(
"pcmpeq.") ||
495 Name.starts_with(
"pcmpgt.") ||
501 Name ==
"pmulhu.w" ||
502 Name ==
"pmulu.dq" ||
503 Name.starts_with(
"pshuf") ||
504 Name.starts_with(
"psll.dq") ||
505 Name.starts_with(
"psrl.dq") ||
506 Name.starts_with(
"psubs.") ||
507 Name.starts_with(
"psubus.") ||
508 Name.starts_with(
"sqrt.p") ||
510 Name ==
"storel.dq" ||
511 Name.starts_with(
"storeu.") ||
514 if (Name.consume_front(
"sse41."))
515 return (Name.starts_with(
"blendp") ||
516 Name ==
"movntdqa" ||
526 Name.starts_with(
"pmovsx") ||
527 Name.starts_with(
"pmovzx") ||
530 if (Name.consume_front(
"sse42."))
531 return Name ==
"crc32.64.8";
533 if (Name.consume_front(
"sse4a."))
534 return Name.starts_with(
"movnt.");
536 if (Name.consume_front(
"ssse3."))
537 return (Name ==
"pabs.b.128" ||
538 Name ==
"pabs.d.128" ||
539 Name ==
"pabs.w.128");
541 if (Name.consume_front(
"xop."))
542 return (Name ==
"vpcmov" ||
543 Name ==
"vpcmov.256" ||
544 Name.starts_with(
"vpcom") ||
545 Name.starts_with(
"vprot"));
547 if (Name.consume_front(
"bmi."))
548 return (Name.starts_with(
"pdep.") ||
549 Name.starts_with(
"pext."));
551 return (Name ==
"addcarry.u32" ||
552 Name ==
"addcarry.u64" ||
553 Name ==
"addcarryx.u32" ||
554 Name ==
"addcarryx.u64" ||
555 Name ==
"subborrow.u32" ||
556 Name ==
"subborrow.u64" ||
557 Name.starts_with(
"vcvtph2ps."));
563 if (!Name.consume_front(
"x86."))
571 if (Name ==
"rdtscp") {
573 if (
F->getFunctionType()->getNumParams() == 0)
578 Intrinsic::x86_rdtscp);
585 if (Name.consume_front(
"sse41.ptest")) {
587 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
588 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
589 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
602 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
603 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
604 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
605 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
606 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
607 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
612 if (Name.consume_front(
"avx512.")) {
613 if (Name.consume_front(
"mask.cmp.")) {
616 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
617 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
618 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
619 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
620 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
621 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
625 }
else if (Name.starts_with(
"vpdpbusd.") ||
626 Name.starts_with(
"vpdpbusds.")) {
629 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
630 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
631 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
632 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
633 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
634 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
638 }
else if (Name.starts_with(
"vpdpwssd.") ||
639 Name.starts_with(
"vpdpwssds.")) {
642 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
643 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
644 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
645 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
646 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
647 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
655 if (Name.consume_front(
"avx2.")) {
656 if (Name.consume_front(
"vpdpb")) {
659 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
660 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
661 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
662 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
663 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
664 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
665 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
666 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
667 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
668 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
669 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
670 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
674 }
else if (Name.consume_front(
"vpdpw")) {
677 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
678 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
679 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
680 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
681 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
682 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
683 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
684 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
685 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
686 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
687 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
688 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
696 if (Name.consume_front(
"avx10.")) {
697 if (Name.consume_front(
"vpdpb")) {
700 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
701 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
702 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
703 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
704 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
705 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
709 }
else if (Name.consume_front(
"vpdpw")) {
711 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
712 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
713 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
714 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
715 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
716 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
724 if (Name.consume_front(
"avx512bf16.")) {
727 .
Case(
"cvtne2ps2bf16.128",
728 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
729 .
Case(
"cvtne2ps2bf16.256",
730 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
731 .
Case(
"cvtne2ps2bf16.512",
732 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
733 .
Case(
"mask.cvtneps2bf16.128",
734 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
735 .
Case(
"cvtneps2bf16.256",
736 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
737 .
Case(
"cvtneps2bf16.512",
738 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
745 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
746 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
747 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
754 if (Name.consume_front(
"xop.")) {
756 if (Name.starts_with(
"vpermil2")) {
759 auto Idx =
F->getFunctionType()->getParamType(2);
760 if (Idx->isFPOrFPVectorTy()) {
761 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
762 unsigned EltSize = Idx->getScalarSizeInBits();
763 if (EltSize == 64 && IdxSize == 128)
764 ID = Intrinsic::x86_xop_vpermil2pd;
765 else if (EltSize == 32 && IdxSize == 128)
766 ID = Intrinsic::x86_xop_vpermil2ps;
767 else if (EltSize == 64 && IdxSize == 256)
768 ID = Intrinsic::x86_xop_vpermil2pd_256;
770 ID = Intrinsic::x86_xop_vpermil2ps_256;
772 }
else if (
F->arg_size() == 2)
775 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
776 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
787 if (Name ==
"seh.recoverfp") {
789 Intrinsic::eh_recoverfp);
801 if (Name.starts_with(
"rbit")) {
804 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
808 if (Name ==
"thread.pointer") {
811 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
815 bool Neon = Name.consume_front(
"neon.");
820 if (Name.consume_front(
"bfdot.")) {
824 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
829 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
830 assert((OperandWidth == 64 || OperandWidth == 128) &&
831 "Unexpected operand width");
833 std::array<Type *, 2> Tys{
844 if (Name.consume_front(
"bfm")) {
846 if (Name.consume_back(
".v4f32.v16i8")) {
892 F->arg_begin()->getType());
896 if (Name.consume_front(
"vst")) {
898 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
902 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
903 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
906 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
907 Intrinsic::arm_neon_vst4lane};
909 auto fArgs =
F->getFunctionType()->params();
910 Type *Tys[] = {fArgs[0], fArgs[1]};
913 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
916 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
925 if (Name.consume_front(
"mve.")) {
927 if (Name ==
"vctp64") {
937 if (Name.starts_with(
"vrintn.v")) {
939 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
944 if (Name.consume_back(
".v4i1")) {
946 if (Name.consume_back(
".predicated.v2i64.v4i32"))
948 return Name ==
"mull.int" || Name ==
"vqdmull";
950 if (Name.consume_back(
".v2i64")) {
952 bool IsGather = Name.consume_front(
"vldr.gather.");
953 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
954 if (Name.consume_front(
"base.")) {
956 Name.consume_front(
"wb.");
959 return Name ==
"predicated.v2i64";
962 if (Name.consume_front(
"offset.predicated."))
963 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
964 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
977 if (Name.consume_front(
"cde.vcx")) {
979 if (Name.consume_back(
".predicated.v2i64.v4i1"))
981 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
982 Name ==
"3q" || Name ==
"3qa";
996 F->arg_begin()->getType());
1002 .
Case(
"smax", Intrinsic::smax)
1003 .
Case(
"smin", Intrinsic::smin)
1004 .
Case(
"umax", Intrinsic::umax)
1005 .
Case(
"umin", Intrinsic::umin)
1008 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1011 F->getReturnType());
1015 if (Name.starts_with(
"addp")) {
1017 if (
F->arg_size() != 2)
1020 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1022 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1028 if (Name.starts_with(
"bfcvt")) {
1034 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1041 if (Name.consume_front(
"sve.")) {
1043 if (Name.consume_front(
"bf")) {
1044 if (Name ==
"mmla") {
1045 Type *Tys[] = {
F->getReturnType(),
1046 std::next(
F->arg_begin())->getType()};
1048 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1051 if (Name.consume_back(
".lane")) {
1055 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1056 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1057 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1069 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1074 if (Name.consume_front(
"convert.from.svbool")) {
1077 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1080 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1085 if (Name.consume_front(
"convert.to.svbool")) {
1088 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1091 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1096 if (Name.consume_front(
"addqv")) {
1098 if (!
F->getReturnType()->isFPOrFPVectorTy())
1101 auto Args =
F->getFunctionType()->params();
1102 Type *Tys[] = {
F->getReturnType(), Args[1]};
1104 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1108 if (Name.consume_front(
"ld")) {
1110 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1111 if (LdRegex.
match(Name)) {
1117 "Expected 2 arguments for ld* intrinsic.");
1118 Type *PtrTy =
F->getArg(1)->getType();
1121 Intrinsic::aarch64_sve_ld2_sret,
1122 Intrinsic::aarch64_sve_ld3_sret,
1123 Intrinsic::aarch64_sve_ld4_sret,
1126 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1132 if (Name.consume_front(
"tuple.")) {
1134 if (Name.starts_with(
"get")) {
1136 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1138 F->getParent(), Intrinsic::vector_extract, Tys);
1142 if (Name.starts_with(
"set")) {
1144 auto Args =
F->getFunctionType()->params();
1145 Type *Tys[] = {Args[0], Args[2], Args[1]};
1147 F->getParent(), Intrinsic::vector_insert, Tys);
1151 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1152 if (CreateTupleRegex.
match(Name)) {
1154 auto Args =
F->getFunctionType()->params();
1155 Type *Tys[] = {
F->getReturnType(), Args[1]};
1157 F->getParent(), Intrinsic::vector_insert, Tys);
1163 if (Name.starts_with(
"rev.nxv")) {
1166 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1172 if (Name.consume_front(
"sme.")) {
1174 if (Name.consume_front(
"ftmopa.")) {
1179 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1180 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1198#define NVVM_TMA_G2S_MODES(M) \
1199 M(tile_1d, "tile.1d") \
1200 M(tile_2d, "tile.2d") \
1201 M(tile_3d, "tile.3d") \
1202 M(tile_4d, "tile.4d") \
1203 M(tile_5d, "tile.5d") \
1204 M(tile_gather4_2d, "tile.gather4.2d") \
1205 M(im2col_3d, "im2col.3d") \
1206 M(im2col_4d, "im2col.4d") \
1207 M(im2col_5d, "im2col.5d") \
1208 M(im2col_w_3d, "im2col.w.3d") \
1209 M(im2col_w_4d, "im2col.w.4d") \
1210 M(im2col_w_5d, "im2col.w.5d") \
1211 M(im2col_w_128_3d, "im2col.w.128.3d") \
1212 M(im2col_w_128_4d, "im2col.w.128.4d") \
1213 M(im2col_w_128_5d, "im2col.w.128.5d")
1225 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1228#define G2S_ID(ID_SUFFIX, NAME) \
1229 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1239 size_t NumParams =
F->getFunctionType()->getNumParams();
1243 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1250 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1251 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1252 "expected the i64 cache-hint after the multicast mask");
1253 Type *MaskTy = Params[MaskIdx];
1268 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1271#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1272 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1284 if (!
F->getFunctionType()
1285 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1301 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1306 size_t NumParams =
F->getFunctionType()->getNumParams();
1307 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1311 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1316 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1329 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1334 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1337 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1357 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1360 auto [RedOpName, ShapeName] = Name.split(
'.');
1365 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1366 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1367 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1368 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1369 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1370 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1371 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1372 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1378 if (Name.consume_front(
"mapa.shared.cluster"))
1379 if (
F->getReturnType()->getPointerAddressSpace() ==
1381 return Intrinsic::nvvm_mapa_shared_cluster;
1383 if (Name.consume_front(
"cp.async.bulk.")) {
1386 .
Case(
"shared.cta.to.cluster",
1387 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1391 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1401 if (!Name.consume_front(
"tcgen05.commit."))
1404 if (Name.consume_front(
"shared."))
1406 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1407 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1410 if (Name.consume_front(
"mc.shared.")) {
1412 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1416 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1417 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1426 if (
F->arg_size() != 2)
1429 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1430 Name.consume_front(
"tcgen05.alloc."))
1432 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1433 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1436 if (Name.consume_front(
"tcgen05.dealloc."))
1438 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1439 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1446 if (Name.consume_front(
"fma.rn."))
1448 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1449 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1450 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1451 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1454 if (Name.consume_front(
"fmax."))
1456 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1457 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1458 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1459 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1460 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1461 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1462 .
Case(
"ftz.nan.xorsign.abs.bf16",
1463 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1464 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1465 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1466 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1467 .
Case(
"ftz.xorsign.abs.bf16x2",
1468 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1469 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1470 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1471 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1472 .
Case(
"nan.xorsign.abs.bf16x2",
1473 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1474 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1475 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1478 if (Name.consume_front(
"fmin."))
1480 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1481 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1482 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1483 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1484 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1485 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1486 .
Case(
"ftz.nan.xorsign.abs.bf16",
1487 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1488 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1489 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1490 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1491 .
Case(
"ftz.xorsign.abs.bf16x2",
1492 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1493 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1494 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1495 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1496 .
Case(
"nan.xorsign.abs.bf16x2",
1497 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1498 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1499 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1502 if (Name.consume_front(
"neg."))
1504 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1505 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1514 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1519 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1522 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1525 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1526 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1534 {Intrinsic::nvvm_fadd, Intrinsic::nvvm_fadd_sat},
1535 {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat}};
1537 {Intrinsic::nvvm_fmul, Intrinsic::nvvm_fmul_sat},
1538 {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat}};
1540static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1542 auto [Modifiers,
Type] = Name.rsplit(
'.');
1544 return std::nullopt;
1554 return std::nullopt;
1556 StringRef Rest = Modifiers.drop_front(2);
1560 return std::nullopt;
1562 return std::make_pair(IIDs[IsFTZ][IsSat], *
RoundingMode);
1566 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1569 return Intrinsic::nvvm_mbarrier_init;
1573 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1574 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1575 Name.consume_front(
"param");
1579 if (!Name.consume_front(
"vp."))
1608 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1609 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1616 if (!Name.consume_front(
"vp."))
1636 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1637 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1642 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1654 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1665 if (Name.starts_with(
"to.fp16")) {
1669 FuncTy->getReturnType());
1672 if (Name.starts_with(
"from.fp16")) {
1676 FuncTy->getReturnType());
1686 if (Defaults.empty())
1689 unsigned FullArgCount = FirstDefault + Defaults.size();
1692 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1695 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1697 NumMissingTrailingParams))
1700 return FullArgCount;
1707 unsigned FullArgCount =
1709 if (FullArgCount == 0)
1715 "total number of default args does not match intrinsic signature");
1720 bool CanUpgradeDebugIntrinsicsToRecords) {
1721 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1726 if (!Name.consume_front(
"llvm.") || Name.empty())
1732 bool IsArm = Name.consume_front(
"arm.");
1733 if (IsArm || Name.consume_front(
"aarch64.")) {
1739 if (Name.consume_front(
"amdgcn.")) {
1740 if (Name ==
"alignbit") {
1743 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1747 if (Name.consume_front(
"atomic.")) {
1748 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1749 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1758 if (Name.starts_with(
"addrspacecast.nonnull")) {
1765 switch (
F->getIntrinsicID()) {
1769 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1770 if (
F->arg_size() == 7) {
1775 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1776 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1777 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1778 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1779 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1780 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1781 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1782 if (
F->arg_size() == 8) {
1789 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1790 Name.consume_front(
"flat.atomic.")) {
1791 if (Name.starts_with(
"fadd") ||
1793 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1794 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1802 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1807 if (Name.starts_with(
"ldexp.")) {
1810 F->getParent(), Intrinsic::ldexp,
1811 {F->getReturnType(), F->getArg(1)->getType()});
1820 if (
F->arg_size() == 1) {
1821 if (Name.consume_front(
"convert.")) {
1835 F->arg_begin()->getType());
1841 if (Name ==
"coro.end" &&
1842 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1843 CoroEndID = Intrinsic::coro_end;
1844 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1845 CoroEndID = Intrinsic::coro_end_async;
1856 if (Name.consume_front(
"dbg.")) {
1858 if (CanUpgradeDebugIntrinsicsToRecords) {
1859 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1860 Name ==
"declare" || Name ==
"label") {
1869 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1872 Intrinsic::dbg_value);
1879 if (Name.consume_front(
"experimental.vector.")) {
1885 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1886 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1887 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1888 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1889 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1891 Intrinsic::vector_partial_reduce_add)
1894 const auto *FT =
F->getFunctionType();
1896 if (ID == Intrinsic::vector_extract ||
1897 ID == Intrinsic::vector_interleave2)
1900 if (ID != Intrinsic::vector_interleave2)
1902 if (ID == Intrinsic::vector_insert ||
1903 ID == Intrinsic::vector_partial_reduce_add)
1911 if (Name.consume_front(
"reduce.")) {
1913 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1914 if (R.match(Name, &
Groups))
1916 .
Case(
"add", Intrinsic::vector_reduce_add)
1917 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1918 .
Case(
"and", Intrinsic::vector_reduce_and)
1919 .
Case(
"or", Intrinsic::vector_reduce_or)
1920 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1921 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1922 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1923 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1924 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1925 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1926 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1931 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1936 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1937 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1942 auto Args =
F->getFunctionType()->params();
1944 {Args[V2 ? 1 : 0]});
1950 if (Name.consume_front(
"splice"))
1954 if (Name.consume_front(
"experimental.stepvector.")) {
1958 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1963 if (Name.starts_with(
"flt.rounds")) {
1966 Intrinsic::get_rounding);
1971 if (Name.starts_with(
"invariant.group.barrier")) {
1973 auto Args =
F->getFunctionType()->params();
1974 Type* ObjectPtr[1] = {Args[0]};
1977 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1982 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1983 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1984 if (IsLifetimeStart || IsLifetimeEnd) {
1985 if (
F->arg_size() == 2) {
1986 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1987 : Intrinsic::lifetime_end;
1992 F->getArg(1)->getType());
1994 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2014 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2015 .StartsWith(
"memmove.", Intrinsic::memmove)
2017 if (
F->arg_size() == 5) {
2021 F->getFunctionType()->params().slice(0, 3);
2027 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2030 const auto *FT =
F->getFunctionType();
2031 Type *ParamTypes[2] = {
2032 FT->getParamType(0),
2036 Intrinsic::memset, ParamTypes);
2042 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2043 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2044 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2045 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2047 if (MaskedID &&
F->arg_size() == 4) {
2049 if (MaskedID == Intrinsic::masked_load ||
2050 MaskedID == Intrinsic::masked_gather) {
2052 F->getParent(), MaskedID,
2053 {F->getReturnType(), F->getArg(0)->getType()});
2057 F->getParent(), MaskedID,
2058 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2064 if (Name.consume_front(
"nvvm.")) {
2066 if (
F->arg_size() == 1) {
2069 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2070 .Case(
"clz.i", Intrinsic::ctlz)
2071 .
Case(
"popc.i", Intrinsic::ctpop)
2075 {F->getReturnType()});
2078 }
else if (
F->arg_size() == 2) {
2081 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2082 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2083 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2084 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2085 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2086 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2090 {F->getReturnType()});
2127 F->getParent(), IID,
F->getReturnType(),
2128 F->getFunctionType()->params());
2139 {F->getArg(0)->getType()});
2187 F->getArg(0)->getType());
2195 bool Expand =
false;
2196 if (Name.consume_front(
"abs."))
2199 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2200 else if (Name.consume_front(
"fabs."))
2202 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2203 else if (Name.consume_front(
"add."))
2206 else if (Name.consume_front(
"mul."))
2209 else if (Name.consume_front(
"ex2.approx."))
2212 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2213 else if (Name.consume_front(
"atomic.load."))
2222 else if (Name.consume_front(
"atomic."))
2237 else if (Name.consume_front(
"bitcast."))
2240 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2241 else if (Name.consume_front(
"rotate."))
2243 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2244 else if (Name.consume_front(
"ptr.gen.to."))
2247 else if (Name.consume_front(
"ptr."))
2250 else if (Name.consume_front(
"ldg.global."))
2252 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2253 Name.starts_with(
"p."));
2256 .
Case(
"barrier0",
true)
2257 .
Case(
"barrier.n",
true)
2258 .
Case(
"barrier.sync.cnt",
true)
2259 .
Case(
"barrier.sync",
true)
2260 .
Case(
"barrier",
true)
2261 .
Case(
"bar.sync",
true)
2262 .
Case(
"barrier0.popc",
true)
2263 .
Case(
"barrier0.and",
true)
2264 .
Case(
"barrier0.or",
true)
2265 .
Case(
"clz.ll",
true)
2266 .
Case(
"popc.ll",
true)
2268 .
Case(
"swap.lo.hi.b64",
true)
2269 .
Case(
"tanh.approx.f32",
true)
2281 if (Name.starts_with(
"objectsize.")) {
2282 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2283 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2286 Intrinsic::objectsize, Tys);
2293 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2296 F->getParent(), Intrinsic::ptr_annotation,
2297 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2303 if (Name.consume_front(
"riscv.")) {
2306 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2307 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2308 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2309 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2312 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2325 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2326 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2335 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2336 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2337 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2338 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2343 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2352 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2354 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2363 if (Name ==
"stackprotectorcheck") {
2367 if (Name.starts_with(
"strip.invariant.group")) {
2372 F->getParent(), Intrinsic::launder_invariant_group,
2373 F->getReturnType());
2379 if (Name ==
"thread.pointer") {
2381 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2387 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2390 F->getParent(), Intrinsic::var_annotation,
2391 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2394 if (Name.consume_front(
"vector.splice")) {
2395 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2405 if (Name.consume_front(
"wasm.")) {
2408 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2409 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2410 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2415 F->getReturnType());
2419 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2421 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2423 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2442 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2452 std::string
Name =
F->getName().str();
2455 Name,
F->getParent());
2466 if (Result != std::nullopt) {
2483 bool CanUpgradeDebugIntrinsicsToRecords) {
2503 GV->
getName() ==
"llvm.global_dtors")) ||
2517 unsigned N =
Init->getNumOperands();
2518 std::vector<Constant *> NewCtors(
N);
2519 for (
unsigned i = 0; i !=
N; ++i) {
2522 Ctor->getAggregateElement(1),
2536 unsigned NumElts = ResultTy->getNumElements() * 8;
2540 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2550 for (
unsigned l = 0; l != NumElts; l += 16)
2551 for (
unsigned i = 0; i != 16; ++i) {
2552 unsigned Idx = NumElts + i - Shift;
2554 Idx -= NumElts - 16;
2555 Idxs[l + i] = Idx + l;
2558 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2562 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2570 unsigned NumElts = ResultTy->getNumElements() * 8;
2574 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2584 for (
unsigned l = 0; l != NumElts; l += 16)
2585 for (
unsigned i = 0; i != 16; ++i) {
2586 unsigned Idx = i + Shift;
2588 Idx += NumElts - 16;
2589 Idxs[l + i] = Idx + l;
2592 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2596 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2604 Mask = Builder.CreateBitCast(Mask, MaskTy);
2610 for (
unsigned i = 0; i != NumElts; ++i)
2612 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2623 if (
C->isAllOnesValue())
2628 return Builder.CreateSelect(Mask, Op0, Op1);
2635 if (
C->isAllOnesValue())
2639 Mask->getType()->getIntegerBitWidth());
2640 Mask = Builder.CreateBitCast(Mask, MaskTy);
2641 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2642 return Builder.CreateSelect(Mask, Op0, Op1);
2655 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2656 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2661 ShiftVal &= (NumElts - 1);
2670 if (ShiftVal > 16) {
2678 for (
unsigned l = 0; l < NumElts; l += 16) {
2679 for (
unsigned i = 0; i != 16; ++i) {
2680 unsigned Idx = ShiftVal + i;
2681 if (!IsVALIGN && Idx >= 16)
2682 Idx += NumElts - 16;
2683 Indices[l + i] = Idx + l;
2688 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2694 bool ZeroMask,
bool IndexForm) {
2697 unsigned EltWidth = Ty->getScalarSizeInBits();
2698 bool IsFloat = Ty->isFPOrFPVectorTy();
2700 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2701 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2702 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2703 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2704 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2705 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2706 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2707 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2708 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2709 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2710 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2711 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2712 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2713 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2714 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2715 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2716 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2717 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2718 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2719 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2720 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2721 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2722 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2723 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2724 else if (VecWidth == 128 && EltWidth == 16)
2725 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2726 else if (VecWidth == 256 && EltWidth == 16)
2727 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2728 else if (VecWidth == 512 && EltWidth == 16)
2729 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2730 else if (VecWidth == 128 && EltWidth == 8)
2731 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2732 else if (VecWidth == 256 && EltWidth == 8)
2733 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2734 else if (VecWidth == 512 && EltWidth == 8)
2735 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2746 Value *V = Builder.CreateIntrinsic(IID, Args);
2758 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2769 bool IsRotateRight) {
2779 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2780 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2783 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2784 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2829 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2834 bool IsShiftRight,
bool ZeroMask) {
2848 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2849 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2852 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2853 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2868 const Align Alignment =
2870 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2875 if (
C->isAllOnesValue())
2876 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2881 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2887 const Align Alignment =
2896 if (
C->isAllOnesValue())
2897 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2902 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2908 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2909 {Op0, Builder.getInt1(
false)});
2924 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2925 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2926 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2927 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2928 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2931 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2932 LHS = Builder.CreateAnd(
LHS, Mask);
2933 RHS = Builder.CreateAnd(
RHS, Mask);
2950 if (!
C || !
C->isAllOnesValue())
2951 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2956 for (
unsigned i = 0; i != NumElts; ++i)
2958 for (
unsigned i = NumElts; i != 8; ++i)
2959 Indices[i] = NumElts + i % NumElts;
2960 Vec = Builder.CreateShuffleVector(Vec,
2964 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2968 unsigned CC,
bool Signed) {
2976 }
else if (CC == 7) {
3012 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3013 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3015 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3016 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3025 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3031 Name = Name.substr(12);
3036 if (Name.starts_with(
"max.p")) {
3037 if (VecWidth == 128 && EltWidth == 32)
3038 IID = Intrinsic::x86_sse_max_ps;
3039 else if (VecWidth == 128 && EltWidth == 64)
3040 IID = Intrinsic::x86_sse2_max_pd;
3041 else if (VecWidth == 256 && EltWidth == 32)
3042 IID = Intrinsic::x86_avx_max_ps_256;
3043 else if (VecWidth == 256 && EltWidth == 64)
3044 IID = Intrinsic::x86_avx_max_pd_256;
3047 }
else if (Name.starts_with(
"min.p")) {
3048 if (VecWidth == 128 && EltWidth == 32)
3049 IID = Intrinsic::x86_sse_min_ps;
3050 else if (VecWidth == 128 && EltWidth == 64)
3051 IID = Intrinsic::x86_sse2_min_pd;
3052 else if (VecWidth == 256 && EltWidth == 32)
3053 IID = Intrinsic::x86_avx_min_ps_256;
3054 else if (VecWidth == 256 && EltWidth == 64)
3055 IID = Intrinsic::x86_avx_min_pd_256;
3058 }
else if (Name.starts_with(
"pshuf.b.")) {
3059 if (VecWidth == 128)
3060 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3061 else if (VecWidth == 256)
3062 IID = Intrinsic::x86_avx2_pshuf_b;
3063 else if (VecWidth == 512)
3064 IID = Intrinsic::x86_avx512_pshuf_b_512;
3067 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3068 if (VecWidth == 128)
3069 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3070 else if (VecWidth == 256)
3071 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3072 else if (VecWidth == 512)
3073 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3076 }
else if (Name.starts_with(
"pmulh.w")) {
3077 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3078 "Unexpected intrinsic");
3081 }
else if (Name.starts_with(
"pmulhu.w")) {
3082 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3083 "Unexpected intrinsic");
3086 }
else if (Name.starts_with(
"pmaddw.d.")) {
3087 if (VecWidth == 128)
3088 IID = Intrinsic::x86_sse2_pmadd_wd;
3089 else if (VecWidth == 256)
3090 IID = Intrinsic::x86_avx2_pmadd_wd;
3091 else if (VecWidth == 512)
3092 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3095 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3096 if (VecWidth == 128)
3097 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3098 else if (VecWidth == 256)
3099 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3100 else if (VecWidth == 512)
3101 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3104 }
else if (Name.starts_with(
"packsswb.")) {
3105 if (VecWidth == 128)
3106 IID = Intrinsic::x86_sse2_packsswb_128;
3107 else if (VecWidth == 256)
3108 IID = Intrinsic::x86_avx2_packsswb;
3109 else if (VecWidth == 512)
3110 IID = Intrinsic::x86_avx512_packsswb_512;
3113 }
else if (Name.starts_with(
"packssdw.")) {
3114 if (VecWidth == 128)
3115 IID = Intrinsic::x86_sse2_packssdw_128;
3116 else if (VecWidth == 256)
3117 IID = Intrinsic::x86_avx2_packssdw;
3118 else if (VecWidth == 512)
3119 IID = Intrinsic::x86_avx512_packssdw_512;
3122 }
else if (Name.starts_with(
"packuswb.")) {
3123 if (VecWidth == 128)
3124 IID = Intrinsic::x86_sse2_packuswb_128;
3125 else if (VecWidth == 256)
3126 IID = Intrinsic::x86_avx2_packuswb;
3127 else if (VecWidth == 512)
3128 IID = Intrinsic::x86_avx512_packuswb_512;
3131 }
else if (Name.starts_with(
"packusdw.")) {
3132 if (VecWidth == 128)
3133 IID = Intrinsic::x86_sse41_packusdw;
3134 else if (VecWidth == 256)
3135 IID = Intrinsic::x86_avx2_packusdw;
3136 else if (VecWidth == 512)
3137 IID = Intrinsic::x86_avx512_packusdw_512;
3140 }
else if (Name.starts_with(
"vpermilvar.")) {
3141 if (VecWidth == 128 && EltWidth == 32)
3142 IID = Intrinsic::x86_avx_vpermilvar_ps;
3143 else if (VecWidth == 128 && EltWidth == 64)
3144 IID = Intrinsic::x86_avx_vpermilvar_pd;
3145 else if (VecWidth == 256 && EltWidth == 32)
3146 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3147 else if (VecWidth == 256 && EltWidth == 64)
3148 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3149 else if (VecWidth == 512 && EltWidth == 32)
3150 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3151 else if (VecWidth == 512 && EltWidth == 64)
3152 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3155 }
else if (Name ==
"cvtpd2dq.256") {
3156 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3157 }
else if (Name ==
"cvtpd2ps.256") {
3158 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3159 }
else if (Name ==
"cvttpd2dq.256") {
3160 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3161 }
else if (Name ==
"cvttps2dq.128") {
3162 IID = Intrinsic::x86_sse2_cvttps2dq;
3163 }
else if (Name ==
"cvttps2dq.256") {
3164 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3165 }
else if (Name.starts_with(
"permvar.")) {
3167 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3168 IID = Intrinsic::x86_avx2_permps;
3169 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3170 IID = Intrinsic::x86_avx2_permd;
3171 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3172 IID = Intrinsic::x86_avx512_permvar_df_256;
3173 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3174 IID = Intrinsic::x86_avx512_permvar_di_256;
3175 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3176 IID = Intrinsic::x86_avx512_permvar_sf_512;
3177 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3178 IID = Intrinsic::x86_avx512_permvar_si_512;
3179 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3180 IID = Intrinsic::x86_avx512_permvar_df_512;
3181 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3182 IID = Intrinsic::x86_avx512_permvar_di_512;
3183 else if (VecWidth == 128 && EltWidth == 16)
3184 IID = Intrinsic::x86_avx512_permvar_hi_128;
3185 else if (VecWidth == 256 && EltWidth == 16)
3186 IID = Intrinsic::x86_avx512_permvar_hi_256;
3187 else if (VecWidth == 512 && EltWidth == 16)
3188 IID = Intrinsic::x86_avx512_permvar_hi_512;
3189 else if (VecWidth == 128 && EltWidth == 8)
3190 IID = Intrinsic::x86_avx512_permvar_qi_128;
3191 else if (VecWidth == 256 && EltWidth == 8)
3192 IID = Intrinsic::x86_avx512_permvar_qi_256;
3193 else if (VecWidth == 512 && EltWidth == 8)
3194 IID = Intrinsic::x86_avx512_permvar_qi_512;
3197 }
else if (Name.starts_with(
"dbpsadbw.")) {
3198 if (VecWidth == 128)
3199 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3200 else if (VecWidth == 256)
3201 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3202 else if (VecWidth == 512)
3203 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3206 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3207 if (VecWidth == 128)
3208 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3209 else if (VecWidth == 256)
3210 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3211 else if (VecWidth == 512)
3212 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3215 }
else if (Name.starts_with(
"conflict.")) {
3216 if (Name[9] ==
'd' && VecWidth == 128)
3217 IID = Intrinsic::x86_avx512_conflict_d_128;
3218 else if (Name[9] ==
'd' && VecWidth == 256)
3219 IID = Intrinsic::x86_avx512_conflict_d_256;
3220 else if (Name[9] ==
'd' && VecWidth == 512)
3221 IID = Intrinsic::x86_avx512_conflict_d_512;
3222 else if (Name[9] ==
'q' && VecWidth == 128)
3223 IID = Intrinsic::x86_avx512_conflict_q_128;
3224 else if (Name[9] ==
'q' && VecWidth == 256)
3225 IID = Intrinsic::x86_avx512_conflict_q_256;
3226 else if (Name[9] ==
'q' && VecWidth == 512)
3227 IID = Intrinsic::x86_avx512_conflict_q_512;
3230 }
else if (Name.starts_with(
"pavg.")) {
3231 if (Name[5] ==
'b' && VecWidth == 128)
3232 IID = Intrinsic::x86_sse2_pavg_b;
3233 else if (Name[5] ==
'b' && VecWidth == 256)
3234 IID = Intrinsic::x86_avx2_pavg_b;
3235 else if (Name[5] ==
'b' && VecWidth == 512)
3236 IID = Intrinsic::x86_avx512_pavg_b_512;
3237 else if (Name[5] ==
'w' && VecWidth == 128)
3238 IID = Intrinsic::x86_sse2_pavg_w;
3239 else if (Name[5] ==
'w' && VecWidth == 256)
3240 IID = Intrinsic::x86_avx2_pavg_w;
3241 else if (Name[5] ==
'w' && VecWidth == 512)
3242 IID = Intrinsic::x86_avx512_pavg_w_512;
3251 Rep = Builder.CreateIntrinsic(IID, Args);
3262 if (AsmStr->find(
"mov\tfp") == 0 &&
3263 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3264 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3265 AsmStr->replace(Pos, 1,
";");
3273 assert(Result &&
"unsupported nvvm.add.*/nvvm.mul.* intrinsic");
3276 return Builder.CreateIntrinsic(
3278 {A, CI->getArgOperand(1),
3279 Builder.getInt32(static_cast<int>(RoundingMode))});
3284 Value *Rep =
nullptr;
3286 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3288 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3289 {Arg, Builder.getTrue()},
3291 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3292 Type *Ty = (Name ==
"abs.bf16")
3296 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3297 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3298 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3299 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3300 : Intrinsic::nvvm_fabs;
3301 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3302 }
else if (Name.consume_front(
"add.")) {
3305 }
else if (Name.consume_front(
"mul.")) {
3308 }
else if (Name.consume_front(
"ex2.approx.")) {
3310 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3311 : Intrinsic::nvvm_ex2_approx;
3312 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3313 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3314 Name.starts_with(
"atomic.load.add.f64.p")) {
3317 Rep = Builder.CreateAtomicRMW(
3323 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3324 Name.starts_with(
"atomic.load.dec.32.p")) {
3329 Rep = Builder.CreateAtomicRMW(
3333 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3339 Op.contains(
".cta.") ?
"block" :
"");
3340 if (
Op.starts_with(
"cas.")) {
3342 Value *Pair = Builder.CreateAtomicCmpXchg(
3345 Rep = Builder.CreateExtractValue(Pair, 0);
3363 "unexpected nvvm scoped atomic intrinsic");
3364 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3367 }
else if (Name ==
"clz.ll") {
3370 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3371 {Arg, Builder.getFalse()},
3373 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3374 }
else if (Name ==
"popc.ll") {
3378 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3379 Arg,
nullptr,
"ctpop");
3380 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3381 }
else if (Name ==
"h2f") {
3383 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3384 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3385 }
else if (Name.consume_front(
"bitcast.") &&
3386 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3389 }
else if (Name ==
"rotate.b32") {
3392 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3393 {Arg, Arg, ShiftAmt});
3394 }
else if (Name ==
"rotate.b64") {
3398 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3399 {Arg, Arg, ZExtShiftAmt});
3400 }
else if (Name ==
"rotate.right.b64") {
3404 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3405 {Arg, Arg, ZExtShiftAmt});
3406 }
else if (Name ==
"swap.lo.hi.b64") {
3409 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3410 {Arg, Arg, Builder.getInt64(32)});
3411 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3414 Name.starts_with(
".to.gen"))) {
3416 }
else if (Name.consume_front(
"ldg.global")) {
3420 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3423 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3425 }
else if (Name ==
"tanh.approx.f32") {
3429 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3431 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3433 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3434 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3436 }
else if (Name ==
"barrier") {
3437 Rep = Builder.CreateIntrinsic(
3438 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3440 }
else if (Name ==
"barrier.sync") {
3441 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3443 }
else if (Name ==
"barrier.sync.cnt") {
3444 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3446 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3447 Name ==
"barrier0.or") {
3449 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3453 .
Case(
"barrier0.popc",
3454 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3455 .
Case(
"barrier0.and",
3456 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3457 .
Case(
"barrier0.or",
3458 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3459 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3460 Rep = Builder.CreateZExt(Bar, CI->
getType());
3474 ? Builder.CreateBitCast(Arg, NewType)
3477 Rep = Builder.CreateCall(NewFn, Args);
3478 if (
F->getReturnType()->isIntegerTy())
3479 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3489 Value *Rep =
nullptr;
3491 if (Name.starts_with(
"sse4a.movnt.")) {
3503 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3506 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3507 }
else if (Name.starts_with(
"avx.movnt.") ||
3508 Name.starts_with(
"avx512.storent.")) {
3520 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3521 }
else if (Name ==
"sse2.storel.dq") {
3526 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3527 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3528 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3529 }
else if (Name.starts_with(
"sse.storeu.") ||
3530 Name.starts_with(
"sse2.storeu.") ||
3531 Name.starts_with(
"avx.storeu.")) {
3534 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3535 }
else if (Name ==
"avx512.mask.store.ss") {
3539 }
else if (Name.starts_with(
"avx512.mask.store")) {
3541 bool Aligned = Name[17] !=
'u';
3544 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3547 bool CmpEq = Name[9] ==
'e';
3550 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3551 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3558 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3559 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3561 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3562 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3563 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3564 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3565 Name.starts_with(
"sse2.sqrt.p") ||
3566 Name.starts_with(
"sse.sqrt.p")) {
3567 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3568 {CI->getArgOperand(0)});
3569 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3573 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3574 : Intrinsic::x86_avx512_sqrt_pd_512;
3577 Rep = Builder.CreateIntrinsic(IID, Args);
3579 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3580 {CI->getArgOperand(0)});
3584 }
else if (Name.starts_with(
"avx512.ptestm") ||
3585 Name.starts_with(
"avx512.ptestnm")) {
3589 Rep = Builder.CreateAnd(Op0, Op1);
3595 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3597 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3600 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3603 }
else if (Name.starts_with(
"avx512.kunpck")) {
3608 for (
unsigned i = 0; i != NumElts; ++i)
3617 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3618 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3619 }
else if (Name ==
"avx512.kand.w") {
3622 Rep = Builder.CreateAnd(
LHS,
RHS);
3623 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3624 }
else if (Name ==
"avx512.kandn.w") {
3627 LHS = Builder.CreateNot(
LHS);
3628 Rep = Builder.CreateAnd(
LHS,
RHS);
3629 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3630 }
else if (Name ==
"avx512.kor.w") {
3633 Rep = Builder.CreateOr(
LHS,
RHS);
3634 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3635 }
else if (Name ==
"avx512.kxor.w") {
3638 Rep = Builder.CreateXor(
LHS,
RHS);
3639 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3640 }
else if (Name ==
"avx512.kxnor.w") {
3643 LHS = Builder.CreateNot(
LHS);
3644 Rep = Builder.CreateXor(
LHS,
RHS);
3645 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3646 }
else if (Name ==
"avx512.knot.w") {
3648 Rep = Builder.CreateNot(Rep);
3649 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3650 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3653 Rep = Builder.CreateOr(
LHS,
RHS);
3654 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3656 if (Name[14] ==
'c')
3660 Rep = Builder.CreateICmpEQ(Rep,
C);
3661 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3662 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3663 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3664 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3665 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3668 ConstantInt::get(I32Ty, 0));
3670 ConstantInt::get(I32Ty, 0));
3672 if (Name.contains(
".add."))
3673 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3674 else if (Name.contains(
".sub."))
3675 EltOp = Builder.CreateFSub(Elt0, Elt1);
3676 else if (Name.contains(
".mul."))
3677 EltOp = Builder.CreateFMul(Elt0, Elt1);
3679 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3680 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3681 ConstantInt::get(I32Ty, 0));
3682 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3684 bool CmpEq = Name[16] ==
'e';
3686 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3688 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3695 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3698 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3701 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3708 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3710 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3711 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3713 if (VecWidth == 128 && EltWidth == 32)
3714 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3715 else if (VecWidth == 256 && EltWidth == 32)
3716 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3717 else if (VecWidth == 512 && EltWidth == 32)
3718 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3719 else if (VecWidth == 128 && EltWidth == 64)
3720 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3721 else if (VecWidth == 256 && EltWidth == 64)
3722 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3723 else if (VecWidth == 512 && EltWidth == 64)
3724 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3731 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3734 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3735 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3737 if (VecWidth == 128 && EltWidth == 32)
3738 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3739 else if (VecWidth == 256 && EltWidth == 32)
3740 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3741 else if (VecWidth == 512 && EltWidth == 32)
3742 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3743 else if (VecWidth == 128 && EltWidth == 64)
3744 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3745 else if (VecWidth == 256 && EltWidth == 64)
3746 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3747 else if (VecWidth == 512 && EltWidth == 64)
3748 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3753 if (VecWidth == 512)
3755 Args.push_back(Mask);
3757 Rep = Builder.CreateIntrinsic(IID, Args);
3758 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3762 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3765 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3766 Name.starts_with(
"avx512.cvtw2mask.") ||
3767 Name.starts_with(
"avx512.cvtd2mask.") ||
3768 Name.starts_with(
"avx512.cvtq2mask.")) {
3773 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3774 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3775 Name.starts_with(
"avx512.mask.pabs")) {
3777 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3778 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3779 Name.starts_with(
"avx512.mask.pmaxs")) {
3781 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3782 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3783 Name.starts_with(
"avx512.mask.pmaxu")) {
3785 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3786 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3787 Name.starts_with(
"avx512.mask.pmins")) {
3789 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3790 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3791 Name.starts_with(
"avx512.mask.pminu")) {
3793 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3794 Name.starts_with(
"avx512.pmulh.w")) {
3796 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3797 Name.starts_with(
"avx512.pmulhu.w")) {
3799 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3800 Name ==
"avx512.pmulu.dq.512" ||
3801 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3803 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3804 Name ==
"avx512.pmul.dq.512" ||
3805 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3807 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3808 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3813 }
else if (Name ==
"avx512.cvtusi2sd") {
3818 }
else if (Name ==
"sse2.cvtss2sd") {
3820 Rep = Builder.CreateFPExt(
3823 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3824 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3825 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3826 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3827 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3828 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3829 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3830 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3831 Name ==
"avx512.mask.cvtqq2ps.256" ||
3832 Name ==
"avx512.mask.cvtqq2ps.512" ||
3833 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3834 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3835 Name ==
"avx.cvt.ps2.pd.256" ||
3836 Name ==
"avx512.mask.cvtps2pd.128" ||
3837 Name ==
"avx512.mask.cvtps2pd.256") {
3842 unsigned NumDstElts = DstTy->getNumElements();
3843 if (NumDstElts < SrcTy->getNumElements()) {
3844 assert(NumDstElts == 2 &&
"Unexpected vector size");
3845 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3848 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3849 bool IsUnsigned = Name.contains(
"cvtu");
3851 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3855 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3856 : Intrinsic::x86_avx512_sitofp_round;
3857 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3860 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3861 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3867 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3868 Name.starts_with(
"vcvtph2ps.")) {
3872 unsigned NumDstElts = DstTy->getNumElements();
3873 if (NumDstElts != SrcTy->getNumElements()) {
3874 assert(NumDstElts == 4 &&
"Unexpected vector size");
3875 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3877 Rep = Builder.CreateBitCast(
3879 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3883 }
else if (Name.starts_with(
"avx512.mask.load")) {
3885 bool Aligned = Name[16] !=
'u';
3888 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3892 ResultTy->getNumElements());
3893 Rep = Builder.CreateIntrinsic(
3894 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3896 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3902 Rep = Builder.CreateIntrinsic(
3903 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3905 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3906 Name.starts_with(
"avx512.mask.expand.")) {
3910 ResultTy->getNumElements());
3912 bool IsCompress = Name[12] ==
'c';
3913 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3914 : Intrinsic::x86_avx512_mask_expand;
3915 Rep = Builder.CreateIntrinsic(
3917 }
else if (Name.starts_with(
"xop.vpcom")) {
3919 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3920 Name.ends_with(
"uq"))
3922 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3923 Name.ends_with(
"d") || Name.ends_with(
"q"))
3932 Name = Name.substr(9);
3933 if (Name.starts_with(
"lt"))
3935 else if (Name.starts_with(
"le"))
3937 else if (Name.starts_with(
"gt"))
3939 else if (Name.starts_with(
"ge"))
3941 else if (Name.starts_with(
"eq"))
3943 else if (Name.starts_with(
"ne"))
3945 else if (Name.starts_with(
"false"))
3947 else if (Name.starts_with(
"true"))
3954 }
else if (Name.starts_with(
"xop.vpcmov")) {
3956 Value *NotSel = Builder.CreateNot(Sel);
3959 Rep = Builder.CreateOr(Sel0, Sel1);
3960 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3961 Name.starts_with(
"avx512.mask.prol")) {
3963 }
else if (Name.starts_with(
"avx512.pror") ||
3964 Name.starts_with(
"avx512.mask.pror")) {
3966 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3967 Name.starts_with(
"avx512.mask.vpshld") ||
3968 Name.starts_with(
"avx512.maskz.vpshld")) {
3969 bool ZeroMask = Name[11] ==
'z';
3971 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3972 Name.starts_with(
"avx512.mask.vpshrd") ||
3973 Name.starts_with(
"avx512.maskz.vpshrd")) {
3974 bool ZeroMask = Name[11] ==
'z';
3976 }
else if (Name ==
"sse42.crc32.64.8") {
3979 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3981 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3982 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3983 Name.starts_with(
"avx512.vbroadcast.s")) {
3986 Type *EltTy = VecTy->getElementType();
3987 unsigned EltNum = VecTy->getNumElements();
3991 for (
unsigned I = 0;
I < EltNum; ++
I)
3992 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3993 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3994 Name.starts_with(
"sse41.pmovzx") ||
3995 Name.starts_with(
"avx2.pmovsx") ||
3996 Name.starts_with(
"avx2.pmovzx") ||
3997 Name.starts_with(
"avx512.mask.pmovsx") ||
3998 Name.starts_with(
"avx512.mask.pmovzx")) {
4000 unsigned NumDstElts = DstTy->getNumElements();
4004 for (
unsigned i = 0; i != NumDstElts; ++i)
4009 bool DoSext = Name.contains(
"pmovsx");
4011 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
4016 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
4017 Name ==
"avx512.mask.pmov.qd.512" ||
4018 Name ==
"avx512.mask.pmov.wb.256" ||
4019 Name ==
"avx512.mask.pmov.wb.512") {
4024 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4025 Name ==
"avx2.vbroadcasti128") {
4031 if (NumSrcElts == 2)
4034 Rep = Builder.CreateShuffleVector(
Load,
4036 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4037 Name.starts_with(
"avx512.mask.shuf.f")) {
4042 unsigned ControlBitsMask = NumLanes - 1;
4043 unsigned NumControlBits = NumLanes / 2;
4046 for (
unsigned l = 0; l != NumLanes; ++l) {
4047 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4049 if (l >= NumLanes / 2)
4050 LaneMask += NumLanes;
4051 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4052 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4058 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4059 Name.starts_with(
"avx512.mask.broadcasti")) {
4062 unsigned NumDstElts =
4066 for (
unsigned i = 0; i != NumDstElts; ++i)
4067 ShuffleMask[i] = i % NumSrcElts;
4073 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4074 Name.starts_with(
"avx2.vbroadcast") ||
4075 Name.starts_with(
"avx512.pbroadcast") ||
4076 Name.starts_with(
"avx512.mask.broadcast.s")) {
4083 Rep = Builder.CreateShuffleVector(
Op, M);
4088 }
else if (Name.starts_with(
"sse2.padds.") ||
4089 Name.starts_with(
"avx2.padds.") ||
4090 Name.starts_with(
"avx512.padds.") ||
4091 Name.starts_with(
"avx512.mask.padds.")) {
4093 }
else if (Name.starts_with(
"sse2.psubs.") ||
4094 Name.starts_with(
"avx2.psubs.") ||
4095 Name.starts_with(
"avx512.psubs.") ||
4096 Name.starts_with(
"avx512.mask.psubs.")) {
4098 }
else if (Name.starts_with(
"sse2.paddus.") ||
4099 Name.starts_with(
"avx2.paddus.") ||
4100 Name.starts_with(
"avx512.mask.paddus.")) {
4102 }
else if (Name.starts_with(
"sse2.psubus.") ||
4103 Name.starts_with(
"avx2.psubus.") ||
4104 Name.starts_with(
"avx512.mask.psubus.")) {
4106 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4111 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4115 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4120 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4125 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4126 Name ==
"avx512.psll.dq.512") {
4130 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4131 Name ==
"avx512.psrl.dq.512") {
4135 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4136 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4137 Name.starts_with(
"avx2.pblendd.")) {
4142 unsigned NumElts = VecTy->getNumElements();
4145 for (
unsigned i = 0; i != NumElts; ++i)
4146 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4148 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4149 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4150 Name ==
"avx2.vinserti128" ||
4151 Name.starts_with(
"avx512.mask.insert")) {
4155 unsigned DstNumElts =
4157 unsigned SrcNumElts =
4159 unsigned Scale = DstNumElts / SrcNumElts;
4166 for (
unsigned i = 0; i != SrcNumElts; ++i)
4168 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4169 Idxs[i] = SrcNumElts;
4170 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4184 for (
unsigned i = 0; i != DstNumElts; ++i)
4187 for (
unsigned i = 0; i != SrcNumElts; ++i)
4188 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4189 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4195 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4196 Name ==
"avx2.vextracti128" ||
4197 Name.starts_with(
"avx512.mask.vextract")) {
4200 unsigned DstNumElts =
4202 unsigned SrcNumElts =
4204 unsigned Scale = SrcNumElts / DstNumElts;
4211 for (
unsigned i = 0; i != DstNumElts; ++i) {
4212 Idxs[i] = i + (
Imm * DstNumElts);
4214 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4220 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4221 Name.starts_with(
"avx512.mask.perm.di.")) {
4225 unsigned NumElts = VecTy->getNumElements();
4228 for (
unsigned i = 0; i != NumElts; ++i)
4229 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4231 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4236 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4248 unsigned HalfSize = NumElts / 2;
4260 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4261 for (
unsigned i = 0; i < HalfSize; ++i)
4262 ShuffleMask[i] = StartIndex + i;
4265 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4266 for (
unsigned i = 0; i < HalfSize; ++i)
4267 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4269 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4271 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4272 Name.starts_with(
"avx512.mask.vpermil.p") ||
4273 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4277 unsigned NumElts = VecTy->getNumElements();
4279 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4280 unsigned IdxMask = ((1 << IdxSize) - 1);
4286 for (
unsigned i = 0; i != NumElts; ++i)
4287 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4289 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4294 }
else if (Name ==
"sse2.pshufl.w" ||
4295 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4300 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4304 for (
unsigned l = 0; l != NumElts; l += 8) {
4305 for (
unsigned i = 0; i != 4; ++i)
4306 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4307 for (
unsigned i = 4; i != 8; ++i)
4308 Idxs[i + l] = i + l;
4311 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4316 }
else if (Name ==
"sse2.pshufh.w" ||
4317 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4322 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4326 for (
unsigned l = 0; l != NumElts; l += 8) {
4327 for (
unsigned i = 0; i != 4; ++i)
4328 Idxs[i + l] = i + l;
4329 for (
unsigned i = 0; i != 4; ++i)
4330 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4333 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4338 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4345 unsigned HalfLaneElts = NumLaneElts / 2;
4348 for (
unsigned i = 0; i != NumElts; ++i) {
4350 Idxs[i] = i - (i % NumLaneElts);
4352 if ((i % NumLaneElts) >= HalfLaneElts)
4356 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4359 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4363 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4364 Name.starts_with(
"avx512.mask.movshdup") ||
4365 Name.starts_with(
"avx512.mask.movsldup")) {
4371 if (Name.starts_with(
"avx512.mask.movshdup."))
4375 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4376 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4377 Idxs[i + l + 0] = i + l +
Offset;
4378 Idxs[i + l + 1] = i + l +
Offset;
4381 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4385 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4386 Name.starts_with(
"avx512.mask.unpckl.")) {
4393 for (
int l = 0; l != NumElts; l += NumLaneElts)
4394 for (
int i = 0; i != NumLaneElts; ++i)
4395 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4397 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4401 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4402 Name.starts_with(
"avx512.mask.unpckh.")) {
4409 for (
int l = 0; l != NumElts; l += NumLaneElts)
4410 for (
int i = 0; i != NumLaneElts; ++i)
4411 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4413 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4417 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4418 Name.starts_with(
"avx512.mask.pand.")) {
4421 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4423 Rep = Builder.CreateBitCast(Rep, FTy);
4426 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4427 Name.starts_with(
"avx512.mask.pandn.")) {
4430 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4431 Rep = Builder.CreateAnd(Rep,
4433 Rep = Builder.CreateBitCast(Rep, FTy);
4436 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4437 Name.starts_with(
"avx512.mask.por.")) {
4440 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4442 Rep = Builder.CreateBitCast(Rep, FTy);
4445 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4446 Name.starts_with(
"avx512.mask.pxor.")) {
4449 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4451 Rep = Builder.CreateBitCast(Rep, FTy);
4454 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4458 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4462 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4466 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4467 if (Name.ends_with(
".512")) {
4469 if (Name[17] ==
's')
4470 IID = Intrinsic::x86_avx512_add_ps_512;
4472 IID = Intrinsic::x86_avx512_add_pd_512;
4474 Rep = Builder.CreateIntrinsic(
4482 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4483 if (Name.ends_with(
".512")) {
4485 if (Name[17] ==
's')
4486 IID = Intrinsic::x86_avx512_div_ps_512;
4488 IID = Intrinsic::x86_avx512_div_pd_512;
4490 Rep = Builder.CreateIntrinsic(
4498 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4499 if (Name.ends_with(
".512")) {
4501 if (Name[17] ==
's')
4502 IID = Intrinsic::x86_avx512_mul_ps_512;
4504 IID = Intrinsic::x86_avx512_mul_pd_512;
4506 Rep = Builder.CreateIntrinsic(
4514 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4515 if (Name.ends_with(
".512")) {
4517 if (Name[17] ==
's')
4518 IID = Intrinsic::x86_avx512_sub_ps_512;
4520 IID = Intrinsic::x86_avx512_sub_pd_512;
4522 Rep = Builder.CreateIntrinsic(
4530 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4531 Name.starts_with(
"avx512.mask.min.p")) &&
4532 Name.drop_front(18) ==
".512") {
4533 bool IsDouble = Name[17] ==
'd';
4534 bool IsMin = Name[13] ==
'i';
4536 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4537 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4540 Rep = Builder.CreateIntrinsic(
4545 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4547 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4548 {CI->getArgOperand(0), Builder.getInt1(false)});
4551 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4552 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4553 bool IsVariable = Name[16] ==
'v';
4554 char Size = Name[16] ==
'.' ? Name[17]
4555 : Name[17] ==
'.' ? Name[18]
4556 : Name[18] ==
'.' ? Name[19]
4560 if (IsVariable && Name[17] !=
'.') {
4561 if (
Size ==
'd' && Name[17] ==
'2')
4562 IID = Intrinsic::x86_avx2_psllv_q;
4563 else if (
Size ==
'd' && Name[17] ==
'4')
4564 IID = Intrinsic::x86_avx2_psllv_q_256;
4565 else if (
Size ==
's' && Name[17] ==
'4')
4566 IID = Intrinsic::x86_avx2_psllv_d;
4567 else if (
Size ==
's' && Name[17] ==
'8')
4568 IID = Intrinsic::x86_avx2_psllv_d_256;
4569 else if (
Size ==
'h' && Name[17] ==
'8')
4570 IID = Intrinsic::x86_avx512_psllv_w_128;
4571 else if (
Size ==
'h' && Name[17] ==
'1')
4572 IID = Intrinsic::x86_avx512_psllv_w_256;
4573 else if (Name[17] ==
'3' && Name[18] ==
'2')
4574 IID = Intrinsic::x86_avx512_psllv_w_512;
4577 }
else if (Name.ends_with(
".128")) {
4579 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4580 : Intrinsic::x86_sse2_psll_d;
4581 else if (
Size ==
'q')
4582 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4583 : Intrinsic::x86_sse2_psll_q;
4584 else if (
Size ==
'w')
4585 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4586 : Intrinsic::x86_sse2_psll_w;
4589 }
else if (Name.ends_with(
".256")) {
4591 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4592 : Intrinsic::x86_avx2_psll_d;
4593 else if (
Size ==
'q')
4594 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4595 : Intrinsic::x86_avx2_psll_q;
4596 else if (
Size ==
'w')
4597 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4598 : Intrinsic::x86_avx2_psll_w;
4603 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4604 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4605 : Intrinsic::x86_avx512_psll_d_512;
4606 else if (
Size ==
'q')
4607 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4608 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4609 : Intrinsic::x86_avx512_psll_q_512;
4610 else if (
Size ==
'w')
4611 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4612 : Intrinsic::x86_avx512_psll_w_512;
4618 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4619 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4620 bool IsVariable = Name[16] ==
'v';
4621 char Size = Name[16] ==
'.' ? Name[17]
4622 : Name[17] ==
'.' ? Name[18]
4623 : Name[18] ==
'.' ? Name[19]
4627 if (IsVariable && Name[17] !=
'.') {
4628 if (
Size ==
'd' && Name[17] ==
'2')
4629 IID = Intrinsic::x86_avx2_psrlv_q;
4630 else if (
Size ==
'd' && Name[17] ==
'4')
4631 IID = Intrinsic::x86_avx2_psrlv_q_256;
4632 else if (
Size ==
's' && Name[17] ==
'4')
4633 IID = Intrinsic::x86_avx2_psrlv_d;
4634 else if (
Size ==
's' && Name[17] ==
'8')
4635 IID = Intrinsic::x86_avx2_psrlv_d_256;
4636 else if (
Size ==
'h' && Name[17] ==
'8')
4637 IID = Intrinsic::x86_avx512_psrlv_w_128;
4638 else if (
Size ==
'h' && Name[17] ==
'1')
4639 IID = Intrinsic::x86_avx512_psrlv_w_256;
4640 else if (Name[17] ==
'3' && Name[18] ==
'2')
4641 IID = Intrinsic::x86_avx512_psrlv_w_512;
4644 }
else if (Name.ends_with(
".128")) {
4646 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4647 : Intrinsic::x86_sse2_psrl_d;
4648 else if (
Size ==
'q')
4649 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4650 : Intrinsic::x86_sse2_psrl_q;
4651 else if (
Size ==
'w')
4652 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4653 : Intrinsic::x86_sse2_psrl_w;
4656 }
else if (Name.ends_with(
".256")) {
4658 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4659 : Intrinsic::x86_avx2_psrl_d;
4660 else if (
Size ==
'q')
4661 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4662 : Intrinsic::x86_avx2_psrl_q;
4663 else if (
Size ==
'w')
4664 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4665 : Intrinsic::x86_avx2_psrl_w;
4670 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4671 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4672 : Intrinsic::x86_avx512_psrl_d_512;
4673 else if (
Size ==
'q')
4674 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4675 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4676 : Intrinsic::x86_avx512_psrl_q_512;
4677 else if (
Size ==
'w')
4678 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4679 : Intrinsic::x86_avx512_psrl_w_512;
4685 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4686 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4687 bool IsVariable = Name[16] ==
'v';
4688 char Size = Name[16] ==
'.' ? Name[17]
4689 : Name[17] ==
'.' ? Name[18]
4690 : Name[18] ==
'.' ? Name[19]
4694 if (IsVariable && Name[17] !=
'.') {
4695 if (
Size ==
's' && Name[17] ==
'4')
4696 IID = Intrinsic::x86_avx2_psrav_d;
4697 else if (
Size ==
's' && Name[17] ==
'8')
4698 IID = Intrinsic::x86_avx2_psrav_d_256;
4699 else if (
Size ==
'h' && Name[17] ==
'8')
4700 IID = Intrinsic::x86_avx512_psrav_w_128;
4701 else if (
Size ==
'h' && Name[17] ==
'1')
4702 IID = Intrinsic::x86_avx512_psrav_w_256;
4703 else if (Name[17] ==
'3' && Name[18] ==
'2')
4704 IID = Intrinsic::x86_avx512_psrav_w_512;
4707 }
else if (Name.ends_with(
".128")) {
4709 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4710 : Intrinsic::x86_sse2_psra_d;
4711 else if (
Size ==
'q')
4712 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4713 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4714 : Intrinsic::x86_avx512_psra_q_128;
4715 else if (
Size ==
'w')
4716 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4717 : Intrinsic::x86_sse2_psra_w;
4720 }
else if (Name.ends_with(
".256")) {
4722 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4723 : Intrinsic::x86_avx2_psra_d;
4724 else if (
Size ==
'q')
4725 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4726 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4727 : Intrinsic::x86_avx512_psra_q_256;
4728 else if (
Size ==
'w')
4729 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4730 : Intrinsic::x86_avx2_psra_w;
4735 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4736 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4737 : Intrinsic::x86_avx512_psra_d_512;
4738 else if (
Size ==
'q')
4739 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4740 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4741 : Intrinsic::x86_avx512_psra_q_512;
4742 else if (
Size ==
'w')
4743 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4744 : Intrinsic::x86_avx512_psra_w_512;
4750 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4752 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4754 }
else if (Name.ends_with(
".movntdqa")) {
4758 LoadInst *LI = Builder.CreateAlignedLoad(
4763 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4764 Name.starts_with(
"fma.vfmsub.") ||
4765 Name.starts_with(
"fma.vfnmadd.") ||
4766 Name.starts_with(
"fma.vfnmsub.")) {
4767 bool NegMul = Name[6] ==
'n';
4768 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4769 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4780 if (NegMul && !IsScalar)
4781 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4782 if (NegMul && IsScalar)
4783 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4785 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4787 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4791 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4799 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4803 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4804 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4805 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4806 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4807 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4808 bool IsMask3 = Name[11] ==
'3';
4809 bool IsMaskZ = Name[11] ==
'z';
4811 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4812 bool NegMul = Name[2] ==
'n';
4813 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4819 if (NegMul && (IsMask3 || IsMaskZ))
4820 A = Builder.CreateFNeg(
A);
4821 if (NegMul && !(IsMask3 || IsMaskZ))
4822 B = Builder.CreateFNeg(
B);
4824 C = Builder.CreateFNeg(
C);
4826 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4827 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4828 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4835 if (Name.back() ==
'd')
4836 IID = Intrinsic::x86_avx512_vfmadd_f64;
4838 IID = Intrinsic::x86_avx512_vfmadd_f32;
4839 Rep = Builder.CreateIntrinsic(IID,
Ops);
4841 Rep = Builder.CreateFMA(
A,
B,
C);
4850 if (NegAcc && IsMask3)
4855 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4857 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4858 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4859 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4860 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4861 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4862 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4863 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4864 bool IsMask3 = Name[11] ==
'3';
4865 bool IsMaskZ = Name[11] ==
'z';
4867 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4868 bool NegMul = Name[2] ==
'n';
4869 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4875 if (NegMul && (IsMask3 || IsMaskZ))
4876 A = Builder.CreateFNeg(
A);
4877 if (NegMul && !(IsMask3 || IsMaskZ))
4878 B = Builder.CreateFNeg(
B);
4880 C = Builder.CreateFNeg(
C);
4887 if (Name[Name.size() - 5] ==
's')
4888 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4890 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4894 Rep = Builder.CreateFMA(
A,
B,
C);
4902 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4906 if (VecWidth == 128 && EltWidth == 32)
4907 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4908 else if (VecWidth == 256 && EltWidth == 32)
4909 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4910 else if (VecWidth == 128 && EltWidth == 64)
4911 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4912 else if (VecWidth == 256 && EltWidth == 64)
4913 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4919 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4920 Rep = Builder.CreateIntrinsic(IID,
Ops);
4921 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4922 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4923 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4924 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4925 bool IsMask3 = Name[11] ==
'3';
4926 bool IsMaskZ = Name[11] ==
'z';
4928 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4929 bool IsSubAdd = Name[3] ==
's';
4933 if (Name[Name.size() - 5] ==
's')
4934 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4936 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4941 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4943 Rep = Builder.CreateIntrinsic(IID,
Ops);
4952 Value *Odd = Builder.CreateCall(FMA,
Ops);
4953 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4954 Value *Even = Builder.CreateCall(FMA,
Ops);
4960 for (
int i = 0; i != NumElts; ++i)
4961 Idxs[i] = i + (i % 2) * NumElts;
4963 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4971 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4972 Name.starts_with(
"avx512.maskz.pternlog.")) {
4973 bool ZeroMask = Name[11] ==
'z';
4977 if (VecWidth == 128 && EltWidth == 32)
4978 IID = Intrinsic::x86_avx512_pternlog_d_128;
4979 else if (VecWidth == 256 && EltWidth == 32)
4980 IID = Intrinsic::x86_avx512_pternlog_d_256;
4981 else if (VecWidth == 512 && EltWidth == 32)
4982 IID = Intrinsic::x86_avx512_pternlog_d_512;
4983 else if (VecWidth == 128 && EltWidth == 64)
4984 IID = Intrinsic::x86_avx512_pternlog_q_128;
4985 else if (VecWidth == 256 && EltWidth == 64)
4986 IID = Intrinsic::x86_avx512_pternlog_q_256;
4987 else if (VecWidth == 512 && EltWidth == 64)
4988 IID = Intrinsic::x86_avx512_pternlog_q_512;
4994 Rep = Builder.CreateIntrinsic(IID, Args);
4998 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4999 Name.starts_with(
"avx512.maskz.vpmadd52")) {
5000 bool ZeroMask = Name[11] ==
'z';
5001 bool High = Name[20] ==
'h' || Name[21] ==
'h';
5004 if (VecWidth == 128 && !
High)
5005 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
5006 else if (VecWidth == 256 && !
High)
5007 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
5008 else if (VecWidth == 512 && !
High)
5009 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
5010 else if (VecWidth == 128 &&
High)
5011 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
5012 else if (VecWidth == 256 &&
High)
5013 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
5014 else if (VecWidth == 512 &&
High)
5015 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5021 Rep = Builder.CreateIntrinsic(IID, Args);
5025 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5026 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5027 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5028 bool ZeroMask = Name[11] ==
'z';
5029 bool IndexForm = Name[17] ==
'i';
5031 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5032 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5033 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5034 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5035 bool ZeroMask = Name[11] ==
'z';
5036 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5039 if (VecWidth == 128 && !IsSaturating)
5040 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5041 else if (VecWidth == 256 && !IsSaturating)
5042 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5043 else if (VecWidth == 512 && !IsSaturating)
5044 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5045 else if (VecWidth == 128 && IsSaturating)
5046 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5047 else if (VecWidth == 256 && IsSaturating)
5048 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5049 else if (VecWidth == 512 && IsSaturating)
5050 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5060 if (Args[1]->
getType()->isVectorTy() &&
5063 ->isIntegerTy(32) &&
5064 Args[2]->
getType()->isVectorTy() &&
5067 ->isIntegerTy(32)) {
5068 Type *NewArgType =
nullptr;
5069 if (VecWidth == 128)
5071 else if (VecWidth == 256)
5073 else if (VecWidth == 512)
5079 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5080 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5083 Rep = Builder.CreateIntrinsic(IID, Args);
5087 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5088 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5089 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5090 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5091 bool ZeroMask = Name[11] ==
'z';
5092 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5095 if (VecWidth == 128 && !IsSaturating)
5096 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5097 else if (VecWidth == 256 && !IsSaturating)
5098 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5099 else if (VecWidth == 512 && !IsSaturating)
5100 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5101 else if (VecWidth == 128 && IsSaturating)
5102 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5103 else if (VecWidth == 256 && IsSaturating)
5104 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5105 else if (VecWidth == 512 && IsSaturating)
5106 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5116 if (Args[1]->
getType()->isVectorTy() &&
5119 ->isIntegerTy(32) &&
5120 Args[2]->
getType()->isVectorTy() &&
5123 ->isIntegerTy(32)) {
5124 Type *NewArgType =
nullptr;
5125 if (VecWidth == 128)
5127 else if (VecWidth == 256)
5129 else if (VecWidth == 512)
5135 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5136 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5139 Rep = Builder.CreateIntrinsic(IID, Args);
5143 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5144 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5145 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5147 if (Name[0] ==
'a' && Name.back() ==
'2')
5148 IID = Intrinsic::x86_addcarry_32;
5149 else if (Name[0] ==
'a' && Name.back() ==
'4')
5150 IID = Intrinsic::x86_addcarry_64;
5151 else if (Name[0] ==
's' && Name.back() ==
'2')
5152 IID = Intrinsic::x86_subborrow_32;
5153 else if (Name[0] ==
's' && Name.back() ==
'4')
5154 IID = Intrinsic::x86_subborrow_64;
5161 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5164 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5167 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5171 }
else if (Name.starts_with(
"avx512.mask.") &&
5174 }
else if (Name.starts_with(
"bmi.pdep.")) {
5176 }
else if (Name.starts_with(
"bmi.pext.")) {
5186 if (Name.starts_with(
"neon.bfcvt")) {
5187 if (Name.starts_with(
"neon.bfcvtn2")) {
5189 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5191 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5192 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5195 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5196 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5198 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5202 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5203 return Builder.CreateShuffleVector(
5206 return Builder.CreateFPTrunc(CI->
getOperand(0),
5209 }
else if (Name.starts_with(
"sve.fcvt")) {
5212 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5213 .
Case(
"sve.fcvtnt.bf16f32",
5214 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5226 if (Args[1]->
getType() != BadPredTy)
5229 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5230 BadPredTy, Args[1]);
5231 Args[1] = Builder.CreateIntrinsic(
5232 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5234 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5238 if (Name ==
"neon.vcvtfp2hf")
5239 return Builder.CreateBitCast(
5240 Builder.CreateFPTrunc(
5244 if (Name ==
"neon.vcvthf2fp")
5245 return Builder.CreateFPExt(
5246 Builder.CreateBitCast(
5256 if (Name ==
"mve.vctp64.old") {
5259 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5262 Value *C1 = Builder.CreateIntrinsic(
5263 Intrinsic::arm_mve_pred_v2i,
5265 return Builder.CreateIntrinsic(
5266 Intrinsic::arm_mve_pred_i2v,
5268 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5269 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5270 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5271 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5273 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5274 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5275 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5276 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5278 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5279 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5280 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5281 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5282 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5283 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5284 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5285 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5286 std::vector<Type *> Tys;
5290 case Intrinsic::arm_mve_mull_int_predicated:
5291 case Intrinsic::arm_mve_vqdmull_predicated:
5292 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5295 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5296 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5297 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5301 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5305 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5309 case Intrinsic::arm_cde_vcx1q_predicated:
5310 case Intrinsic::arm_cde_vcx1qa_predicated:
5311 case Intrinsic::arm_cde_vcx2q_predicated:
5312 case Intrinsic::arm_cde_vcx2qa_predicated:
5313 case Intrinsic::arm_cde_vcx3q_predicated:
5314 case Intrinsic::arm_cde_vcx3qa_predicated:
5321 std::vector<Value *>
Ops;
5323 Type *Ty =
Op->getType();
5324 if (Ty->getScalarSizeInBits() == 1) {
5325 Value *C1 = Builder.CreateIntrinsic(
5326 Intrinsic::arm_mve_pred_v2i,
5328 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5333 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5348 auto UpgradeLegacyWMMAIUIntrinsicCall =
5353 Args.push_back(Builder.getFalse());
5357 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5364 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5368 NewCall->copyMetadata(*CI);
5372 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5373 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5374 "intrinsic should have 7 arguments");
5377 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5379 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5380 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5381 "intrinsic should have 8 arguments");
5386 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5389 switch (
F->getIntrinsicID()) {
5392 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5393 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5394 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5395 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5396 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5397 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5412 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5415 F->getParent(),
F->getIntrinsicID(), Overloads);
5420 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5424 NewCall->copyMetadata(*CI);
5425 NewCall->takeName(CI);
5430 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5436 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5437 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5445 if (Name.starts_with(
"addrspacecast.nonnull")) {
5448 Value *ASC = Builder.CreateAddrSpaceCast(
5471 if (NumOperands < 3)
5484 bool IsVolatile =
false;
5488 if (NumOperands > 3)
5493 if (NumOperands > 5) {
5495 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5509 if (VT->getElementType()->isIntegerTy(16)) {
5512 Val = Builder.CreateBitCast(Val, AsBF16);
5520 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5522 unsigned AddrSpace = PtrTy->getAddressSpace();
5525 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5527 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5532 MDNode *RangeNotPrivate =
5535 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5541 return Builder.CreateBitCast(RMW, RetTy);
5562 return MAV->getMetadata();
5571 if (Name ==
"label") {
5573 }
else if (Name ==
"assign") {
5580 }
else if (Name ==
"declare") {
5584 }
else if (Name ==
"addr") {
5594 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5595 }
else if (Name ==
"value") {
5598 unsigned ExprOp = 2;
5613 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5621 int64_t OffsetVal =
Offset->getSExtValue();
5622 return Builder.CreateIntrinsic(OffsetVal >= 0
5623 ? Intrinsic::vector_splice_left
5624 : Intrinsic::vector_splice_right,
5626 {CI->getArgOperand(0), CI->getArgOperand(1),
5627 Builder.getInt32(std::abs(OffsetVal))});
5632 if (Name.starts_with(
"to.fp16")) {
5634 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5635 return Builder.CreateBitCast(Cast, CI->
getType());
5638 if (Name.starts_with(
"from.fp16")) {
5640 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5641 return Builder.CreateFPExt(Cast, CI->
getType());
5700 else if (Opcode == Instruction::ICmp)
5703 else if (Opcode == Instruction::FCmp)
5706 else if (Opcode == Instruction::Select)
5711 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5723 if (Defaults.empty())
5726 unsigned OldArgCount = CI->
arg_size();
5727 unsigned NewArgCount = NewFn->
arg_size();
5729 if (OldArgCount < FirstDefault)
5733 if (OldArgCount > NewArgCount)
5738 if (OldArgCount == NewArgCount) {
5750 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5751 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5752 "missing argument outside the default range");
5753 Type *ParamTy = NewFT->getParamType(Idx);
5758 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5764 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5795 if (!Name.consume_front(
"llvm."))
5798 bool IsX86 = Name.consume_front(
"x86.");
5799 bool IsNVVM = Name.consume_front(
"nvvm.");
5800 bool IsAArch64 = Name.consume_front(
"aarch64.");
5801 bool IsARM = Name.consume_front(
"arm.");
5802 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5803 bool IsDbg = Name.consume_front(
"dbg.");
5805 (Name.consume_front(
"experimental.vector.splice") ||
5806 Name.consume_front(
"vector.splice")) &&
5807 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5808 Value *Rep =
nullptr;
5810 if (!IsX86 && Name ==
"stackprotectorcheck") {
5812 }
else if (IsNVVM) {
5816 }
else if (IsAArch64) {
5820 }
else if (IsAMDGCN) {
5824 }
else if (IsOldSplice) {
5826 }
else if (Name.consume_front(
"convert.")) {
5828 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5843 const auto &DefaultCase = [&]() ->
void {
5851 "Unknown function for CallBase upgrade and isn't just a name change");
5859 "Return type must have changed");
5860 assert(OldST->getNumElements() ==
5862 "Must have same number of elements");
5865 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5868 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5869 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5870 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5891 case Intrinsic::arm_neon_vst1:
5892 case Intrinsic::arm_neon_vst2:
5893 case Intrinsic::arm_neon_vst3:
5894 case Intrinsic::arm_neon_vst4:
5895 case Intrinsic::arm_neon_vst2lane:
5896 case Intrinsic::arm_neon_vst3lane:
5897 case Intrinsic::arm_neon_vst4lane: {
5899 NewCall = Builder.CreateCall(NewFn, Args);
5902 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5903 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5904 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5909 NewCall = Builder.CreateCall(NewFn, Args);
5912 case Intrinsic::aarch64_sve_ld3_sret:
5913 case Intrinsic::aarch64_sve_ld4_sret:
5914 case Intrinsic::aarch64_sve_ld2_sret: {
5922 Name = Name.substr(5);
5929 unsigned MinElts = RetTy->getMinNumElements() /
N;
5931 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5933 for (
unsigned I = 0;
I <
N;
I++) {
5934 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5935 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5941 case Intrinsic::coro_end_async:
5942 case Intrinsic::coro_end: {
5944 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5946 NewCall = Builder.CreateCall(NewFn, Args);
5951 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5952 Value *InRamp = Builder.CreateCall(IsInRamp);
5962 case Intrinsic::vector_extract: {
5964 Name = Name.substr(5);
5965 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5970 unsigned MinElts = RetTy->getMinNumElements();
5973 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5977 case Intrinsic::vector_insert: {
5979 Name = Name.substr(5);
5980 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5984 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5989 NewCall = Builder.CreateCall(
5993 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5999 assert(
N > 1 &&
"Create is expected to be between 2-4");
6002 unsigned MinElts = RetTy->getMinNumElements() /
N;
6003 for (
unsigned I = 0;
I <
N;
I++) {
6005 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
6012 case Intrinsic::arm_neon_bfdot:
6013 case Intrinsic::arm_neon_bfmmla:
6014 case Intrinsic::arm_neon_bfmlalb:
6015 case Intrinsic::arm_neon_bfmlalt:
6016 case Intrinsic::aarch64_neon_bfdot:
6017 case Intrinsic::aarch64_neon_bfmmla:
6018 case Intrinsic::aarch64_neon_bfmlalb:
6019 case Intrinsic::aarch64_neon_bfmlalt: {
6022 "Mismatch between function args and call args");
6023 size_t OperandWidth =
6025 assert((OperandWidth == 64 || OperandWidth == 128) &&
6026 "Unexpected operand width");
6028 auto Iter = CI->
args().begin();
6029 Args.push_back(*Iter++);
6030 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6031 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6032 NewCall = Builder.CreateCall(NewFn, Args);
6036 case Intrinsic::bitreverse:
6037 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6040 case Intrinsic::ctlz:
6041 case Intrinsic::cttz: {
6048 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6052 case Intrinsic::objectsize: {
6053 Value *NullIsUnknownSize =
6057 NewCall = Builder.CreateCall(
6062 case Intrinsic::ctpop:
6063 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6065 case Intrinsic::dbg_value: {
6067 Name = Name.substr(5);
6069 if (Name.starts_with(
"dbg.addr")) {
6083 if (
Offset->isNullValue()) {
6084 NewCall = Builder.CreateCall(
6093 case Intrinsic::ptr_annotation:
6101 NewCall = Builder.CreateCall(
6110 case Intrinsic::var_annotation:
6117 NewCall = Builder.CreateCall(
6126 case Intrinsic::riscv_aes32dsi:
6127 case Intrinsic::riscv_aes32dsmi:
6128 case Intrinsic::riscv_aes32esi:
6129 case Intrinsic::riscv_aes32esmi:
6130 case Intrinsic::riscv_sm4ks:
6131 case Intrinsic::riscv_sm4ed: {
6141 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6142 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6148 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6149 Value *Res = NewCall;
6151 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6157 case Intrinsic::nvvm_mapa_shared_cluster: {
6161 Value *Res = NewCall;
6162 Res = Builder.CreateAddrSpaceCast(
6169 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6171 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6173 Args[0] = Builder.CreateAddrSpaceCast(
6177 Args.push_back(Builder.getInt32(0));
6179 NewCall = Builder.CreateCall(NewFn, Args);
6185 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6190 for (
unsigned I = 0;
I < 4; ++
I)
6192 Args.push_back(Builder.getInt32(0));
6193 Args.push_back(Builder.getInt32(0));
6196 Args.push_back(Builder.getInt1(
false));
6197 Args.push_back(Builder.getInt32(0));
6199 NewCall = Builder.CreateCall(NewFn, Args);
6205 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6208 Args[0] = Builder.CreateAddrSpaceCast(
6211 NewCall = Builder.CreateCall(NewFn, Args);
6218#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6219 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6221#undef G2S_CLUSTER_CASE
6226 Args[0] = Builder.CreateAddrSpaceCast(
6232 Args.push_back(Builder.getInt32(0));
6234 NewCall = Builder.CreateCall(NewFn, Args);
6241#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6242 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6250 "expected only the trailing flag_valid_pattern to be missing");
6251 Args.push_back(Builder.getInt32(0));
6253 NewCall = Builder.CreateCall(NewFn, Args);
6259#undef NVVM_TMA_G2S_MODES
6262 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6263 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6264 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6265 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6266 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6267 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6268 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6269 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6271 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6275 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6276 NewCall = Builder.CreateCall(NewFn, Args);
6279 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6280 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6281 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6282 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6285 Builder.getFalse()});
6287 case Intrinsic::nvvm_mbarrier_init: {
6291 if (Args.size() == 2)
6292 Args.push_back(Builder.getInt32(0));
6293 NewCall = Builder.CreateCall(NewFn, Args);
6296 case Intrinsic::riscv_sha256sig0:
6297 case Intrinsic::riscv_sha256sig1:
6298 case Intrinsic::riscv_sha256sum0:
6299 case Intrinsic::riscv_sha256sum1:
6300 case Intrinsic::riscv_sm3p0:
6301 case Intrinsic::riscv_sm3p1: {
6308 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6310 NewCall = Builder.CreateCall(NewFn, Arg);
6312 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6319 case Intrinsic::x86_xop_vfrcz_ss:
6320 case Intrinsic::x86_xop_vfrcz_sd:
6321 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6324 case Intrinsic::x86_xop_vpermil2pd:
6325 case Intrinsic::x86_xop_vpermil2ps:
6326 case Intrinsic::x86_xop_vpermil2pd_256:
6327 case Intrinsic::x86_xop_vpermil2ps_256: {
6331 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6332 NewCall = Builder.CreateCall(NewFn, Args);
6336 case Intrinsic::x86_sse41_ptestc:
6337 case Intrinsic::x86_sse41_ptestz:
6338 case Intrinsic::x86_sse41_ptestnzc: {
6352 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6353 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6355 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6359 case Intrinsic::x86_rdtscp: {
6365 NewCall = Builder.CreateCall(NewFn);
6367 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6370 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6378 case Intrinsic::x86_sse41_insertps:
6379 case Intrinsic::x86_sse41_dppd:
6380 case Intrinsic::x86_sse41_dpps:
6381 case Intrinsic::x86_sse41_mpsadbw:
6382 case Intrinsic::x86_avx_dp_ps_256:
6383 case Intrinsic::x86_avx2_mpsadbw: {
6389 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6390 NewCall = Builder.CreateCall(NewFn, Args);
6394 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6395 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6396 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6397 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6398 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6399 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6405 NewCall = Builder.CreateCall(NewFn, Args);
6414 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6415 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6416 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6417 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6418 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6419 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6423 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6424 Args[1] = Builder.CreateBitCast(
6427 NewCall = Builder.CreateCall(NewFn, Args);
6428 Value *Res = Builder.CreateBitCast(
6436 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6437 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6438 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6442 Args[1] = Builder.CreateBitCast(
6444 Args[2] = Builder.CreateBitCast(
6447 NewCall = Builder.CreateCall(NewFn, Args);
6451 case Intrinsic::thread_pointer: {
6452 NewCall = Builder.CreateCall(NewFn, {});
6456 case Intrinsic::memcpy:
6457 case Intrinsic::memmove:
6458 case Intrinsic::memset: {
6474 NewCall = Builder.CreateCall(NewFn, Args);
6477 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6478 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6479 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6484 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6487 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6491 case Intrinsic::masked_load:
6492 case Intrinsic::masked_gather:
6493 case Intrinsic::masked_store:
6494 case Intrinsic::masked_scatter: {
6500 auto GetMaybeAlign = [](
Value *
Op) {
6502 uint64_t Val = CI->getZExtValue();
6510 auto GetAlign = [&](
Value *
Op) {
6519 case Intrinsic::masked_load:
6520 NewCall = Builder.CreateMaskedLoad(
6524 case Intrinsic::masked_gather:
6525 NewCall = Builder.CreateMaskedGather(
6531 case Intrinsic::masked_store:
6532 NewCall = Builder.CreateMaskedStore(
6536 case Intrinsic::masked_scatter:
6537 NewCall = Builder.CreateMaskedScatter(
6539 DL.getValueOrABITypeAlignment(
6553 case Intrinsic::lifetime_start:
6554 case Intrinsic::lifetime_end: {
6566 NewCall = Builder.CreateLifetimeStart(Ptr);
6568 NewCall = Builder.CreateLifetimeEnd(Ptr);
6577 case Intrinsic::x86_avx512_vpdpbusd_128:
6578 case Intrinsic::x86_avx512_vpdpbusd_256:
6579 case Intrinsic::x86_avx512_vpdpbusd_512:
6580 case Intrinsic::x86_avx512_vpdpbusds_128:
6581 case Intrinsic::x86_avx512_vpdpbusds_256:
6582 case Intrinsic::x86_avx512_vpdpbusds_512:
6583 case Intrinsic::x86_avx2_vpdpbssd_128:
6584 case Intrinsic::x86_avx2_vpdpbssd_256:
6585 case Intrinsic::x86_avx10_vpdpbssd_512:
6586 case Intrinsic::x86_avx2_vpdpbssds_128:
6587 case Intrinsic::x86_avx2_vpdpbssds_256:
6588 case Intrinsic::x86_avx10_vpdpbssds_512:
6589 case Intrinsic::x86_avx2_vpdpbsud_128:
6590 case Intrinsic::x86_avx2_vpdpbsud_256:
6591 case Intrinsic::x86_avx10_vpdpbsud_512:
6592 case Intrinsic::x86_avx2_vpdpbsuds_128:
6593 case Intrinsic::x86_avx2_vpdpbsuds_256:
6594 case Intrinsic::x86_avx10_vpdpbsuds_512:
6595 case Intrinsic::x86_avx2_vpdpbuud_128:
6596 case Intrinsic::x86_avx2_vpdpbuud_256:
6597 case Intrinsic::x86_avx10_vpdpbuud_512:
6598 case Intrinsic::x86_avx2_vpdpbuuds_128:
6599 case Intrinsic::x86_avx2_vpdpbuuds_256:
6600 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6605 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6606 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6608 NewCall = Builder.CreateCall(NewFn, Args);
6611 case Intrinsic::x86_avx512_vpdpwssd_128:
6612 case Intrinsic::x86_avx512_vpdpwssd_256:
6613 case Intrinsic::x86_avx512_vpdpwssd_512:
6614 case Intrinsic::x86_avx512_vpdpwssds_128:
6615 case Intrinsic::x86_avx512_vpdpwssds_256:
6616 case Intrinsic::x86_avx512_vpdpwssds_512:
6617 case Intrinsic::x86_avx2_vpdpwsud_128:
6618 case Intrinsic::x86_avx2_vpdpwsud_256:
6619 case Intrinsic::x86_avx10_vpdpwsud_512:
6620 case Intrinsic::x86_avx2_vpdpwsuds_128:
6621 case Intrinsic::x86_avx2_vpdpwsuds_256:
6622 case Intrinsic::x86_avx10_vpdpwsuds_512:
6623 case Intrinsic::x86_avx2_vpdpwusd_128:
6624 case Intrinsic::x86_avx2_vpdpwusd_256:
6625 case Intrinsic::x86_avx10_vpdpwusd_512:
6626 case Intrinsic::x86_avx2_vpdpwusds_128:
6627 case Intrinsic::x86_avx2_vpdpwusds_256:
6628 case Intrinsic::x86_avx10_vpdpwusds_512:
6629 case Intrinsic::x86_avx2_vpdpwuud_128:
6630 case Intrinsic::x86_avx2_vpdpwuud_256:
6631 case Intrinsic::x86_avx10_vpdpwuud_512:
6632 case Intrinsic::x86_avx2_vpdpwuuds_128:
6633 case Intrinsic::x86_avx2_vpdpwuuds_256:
6634 case Intrinsic::x86_avx10_vpdpwuuds_512:
6639 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6640 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6642 NewCall = Builder.CreateCall(NewFn, Args);
6645 assert(NewCall &&
"Should have either set this variable or returned through "
6646 "the default case");
6653 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6667 F->eraseFromParent();
6673 if (NumOperands == 0)
6681 if (NumOperands == 3) {
6685 Metadata *Elts2[] = {ScalarType, ScalarType,
6701 if (NumOperands == 0 || NumOperands % 3 != 0)
6706 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6711 if (Upgraded ==
Tag)
6721 if (
Opc != Instruction::BitCast)
6725 Type *SrcTy = V->getType();
6742 if (
Opc != Instruction::BitCast)
6745 Type *SrcTy =
C->getType();
6762 if (Flag.getNumOperands() < 3)
6763 return std::nullopt;
6765 return Name->getString();
6766 return std::nullopt;
6780 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6781 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6782 if (auto Name = getModuleFlagNameSafely(*Flag))
6783 return *Name ==
"Debug Info Version";
6786 if (OpIt != ModFlags->op_end()) {
6787 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6794 bool BrokenDebugInfo =
false;
6797 if (!BrokenDebugInfo)
6803 M.getContext().diagnose(Diag);
6810 M.getContext().diagnose(DiagVersion);
6820 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6823 if (
F->hasFnAttribute(Attr)) {
6826 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6828 auto [Part, Rest] = S.
split(
',');
6834 const unsigned Dim = DimC -
'x';
6835 assert(Dim < 3 &&
"Unexpected dim char");
6845 F->addFnAttr(Attr, NewAttr);
6849 return S ==
"x" || S ==
"y" || S ==
"z";
6854 if (
K ==
"kernel") {
6866 const unsigned Idx = (AlignIdxValuePair >> 16);
6867 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6872 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6877 if (
K ==
"minctasm") {
6882 if (
K ==
"maxnreg") {
6887 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6891 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6895 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6899 if (
K ==
"grid_constant") {
6914 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6921 if (!SeenNodes.
insert(MD).second)
6928 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6935 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6937 const MDOperand &V = MD->getOperand(j + 1);
6943 if (NewOperands.
size() > 1)
6956 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6957 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6958 if (ModRetainReleaseMarker) {
6964 ID->getString().split(ValueComp,
"#");
6965 if (ValueComp.
size() == 2) {
6966 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6970 M.eraseNamedMetadata(ModRetainReleaseMarker);
6981 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
7007 bool InvalidCast =
false;
7009 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7022 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7024 Args.push_back(Arg);
7031 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7036 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7049 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7057 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7058 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7059 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7060 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7061 {
"objc_autoreleaseReturnValue",
7062 llvm::Intrinsic::objc_autoreleaseReturnValue},
7063 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7064 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7065 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7066 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7067 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7068 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7069 {
"objc_release", llvm::Intrinsic::objc_release},
7070 {
"objc_retain", llvm::Intrinsic::objc_retain},
7071 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7072 {
"objc_retainAutoreleaseReturnValue",
7073 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7074 {
"objc_retainAutoreleasedReturnValue",
7075 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7076 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7077 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7078 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7079 {
"objc_unsafeClaimAutoreleasedReturnValue",
7080 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7081 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7082 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7083 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7084 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7085 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7086 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7087 {
"objc_arc_annotation_topdown_bbstart",
7088 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7089 {
"objc_arc_annotation_topdown_bbend",
7090 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7091 {
"objc_arc_annotation_bottomup_bbstart",
7092 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7093 {
"objc_arc_annotation_bottomup_bbend",
7094 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7096 for (
auto &
I : RuntimeFuncs)
7097 UpgradeToIntrinsic(
I.first,
I.second);
7121 std::optional<bool> UseAddressDisc;
7124 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7125 for (
const MDNode *Flag : ModFlags->operands()) {
7127 if (Name && (*Name ==
"ptrauth-init-fini" ||
7128 *Name ==
"ptrauth-init-fini-address-discrimination"))
7133 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7134 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7135 constexpr unsigned ExpectedAddressMarker = 1;
7138 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7141 bool HasAddressDisc;
7142 if (!CPA->hasAddressDiscriminator())
7143 HasAddressDisc =
false;
7144 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7145 HasAddressDisc =
true;
7149 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7152 UseAddressDisc = HasAddressDisc;
7153 return CPA->getPointer();
7157 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7160 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7162 if (!GV || !GV->hasInitializer())
7166 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7169 std::vector<Constant *> NewStructors;
7170 NewStructors.reserve(OldStructorsArray->getNumOperands());
7172 for (
Use &U : OldStructorsArray->operands()) {
7181 Func = UpgradeSinglePointer(Func);
7185 NewStructors.push_back(
7194 if (GlobalArraysToUpgrade.
empty())
7196 assert(UseAddressDisc.has_value());
7198 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7199 GV->setInitializer(NewInit);
7202 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7212 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7216 bool HasObjCFlag =
false, HasClassProperties =
false;
7217 bool HasSwiftVersionFlag =
false;
7218 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7225 if (
Op->getNumOperands() != 3)
7239 if (ID->getString() ==
"Objective-C Image Info Version")
7241 if (ID->getString() ==
"Objective-C Class Properties")
7242 HasClassProperties =
true;
7244 if (ID->getString() ==
"PIC Level") {
7245 if (
auto *Behavior =
7247 uint64_t V = Behavior->getLimitedValue();
7253 if (ID->getString() ==
"PIE Level")
7254 if (
auto *Behavior =
7262 if (ID->getString() ==
"branch-target-enforcement" ||
7263 (ID->getString().starts_with(
"sign-return-address") &&
7264 ID->getString() !=
"sign-return-address-harden")) {
7265 if (
auto *Behavior =
7271 Op->getOperand(1),
Op->getOperand(2)};
7281 if (ID->getString() ==
"Objective-C Image Info Section") {
7284 Value->getString().split(ValueComp,
" ");
7285 if (ValueComp.
size() != 1) {
7286 std::string NewValue;
7287 for (
auto &S : ValueComp)
7288 NewValue += S.str();
7299 if (ID->getString() ==
"Objective-C Garbage Collection") {
7302 assert(Md->getValue() &&
"Expected non-empty metadata");
7303 auto Type = Md->getValue()->getType();
7306 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7307 if ((Val & 0xff) != Val) {
7308 HasSwiftVersionFlag =
true;
7309 SwiftABIVersion = (Val & 0xff00) >> 8;
7310 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7311 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7322 if (ID->getString() ==
"amdgpu_code_object_version") {
7325 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7334 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7363 if (HasObjCFlag && !HasClassProperties) {
7369 if (HasSwiftVersionFlag) {
7373 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7375 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7383 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7387 auto MatchesVersion = [](
const MDNode *
Op) {
7388 return Op->getNumOperands() >= 3 &&
7402 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7403 assert(
Op->getNumOperands() >= 2 &&
7404 "Expected at least 2 operands - name and linkage type");
7416 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7427 auto TrimSpaces = [](
StringRef Section) -> std::string {
7429 Section.split(Components,
',');
7434 for (
auto Component : Components)
7435 OS <<
',' << Component.trim();
7440 for (
auto &GV : M.globals()) {
7441 if (!GV.hasSection())
7446 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7451 GV.setSection(TrimSpaces(Section));
7467struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7468 StrictFPUpgradeVisitor() =
default;
7471 if (!
Call.isStrictFP())
7477 Call.removeFnAttr(Attribute::StrictFP);
7478 Call.addFnAttr(Attribute::NoBuiltin);
7483struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7484 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7485 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7487 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7502 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7503 StrictFPUpgradeVisitor SFPV;
7508 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7509 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7510 for (
auto &Arg :
F.args())
7512 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7514 bool AddingAttrs =
false, RemovingAttrs =
false;
7515 AttrBuilder AttrsToAdd(
F.getContext());
7520 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7521 A.isValid() &&
A.isStringAttribute()) {
7522 F.setSection(
A.getValueAsString());
7524 RemovingAttrs =
true;
7528 A.isValid() &&
A.isStringAttribute()) {
7531 AddingAttrs = RemovingAttrs =
true;
7534 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7535 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7537 RemovingAttrs =
true;
7538 if (
A.getValueAsString() ==
"true") {
7539 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7548 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7551 if (
A.getValueAsBool()) {
7552 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7558 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7559 RemovingAttrs =
true;
7566 bool HandleDenormalMode =
false;
7568 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7571 DenormalFPMath = ParsedMode;
7573 AddingAttrs = RemovingAttrs =
true;
7574 HandleDenormalMode =
true;
7578 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7582 DenormalFPMathF32 = ParsedMode;
7584 AddingAttrs = RemovingAttrs =
true;
7585 HandleDenormalMode =
true;
7589 if (HandleDenormalMode)
7590 AttrsToAdd.addDenormalFPEnvAttr(
7594 F.removeFnAttrs(AttrsToRemove);
7597 F.addFnAttrs(AttrsToAdd);
7603 if (!
F.hasFnAttribute(FnAttrName)) {
7604 F.addFnAttr(FnAttrName,
Value);
7606 <<
"\", function: " <<
F.getName() <<
"\n");
7614 if (!
F.hasFnAttribute(FnAttrName)) {
7616 F.addFnAttr(FnAttrName);
7618 <<
", function: " <<
F.getName() <<
"\n");
7621 auto A =
F.getFnAttribute(FnAttrName);
7622 if (
"false" ==
A.getValueAsString()) {
7623 F.removeFnAttr(FnAttrName);
7625 <<
"=\"false\", function: " <<
F.getName() <<
"\n");
7626 }
else if (
"true" ==
A.getValueAsString()) {
7627 F.removeFnAttr(FnAttrName);
7628 F.addFnAttr(FnAttrName);
7630 <<
"=\"true\", function: " <<
F.getName() <<
"\n");
7637 M.setModuleFlag(Behavior,
Key, Val);
7638 LLVM_DEBUG(
dbgs() <<
"Converted module flag: " <<
"{" << Behavior <<
", "
7639 <<
Key <<
", " << Val <<
"}\n");
7643 Triple T(M.getTargetTriple());
7644 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7647 uint64_t BTEValue = 0;
7648 uint64_t BPPLRValue = 0;
7649 uint64_t GCSValue = 0;
7650 uint64_t SRAValue = 0;
7651 uint64_t SRAALLValue = 0;
7652 uint64_t SRABKeyValue = 0;
7654 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7658 if (
Op->getNumOperands() != 3)
7667 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7668 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7669 : IDStr ==
"guarded-control-stack" ? &GCSValue
7670 : IDStr ==
"sign-return-address" ? &SRAValue
7671 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7672 : IDStr ==
"sign-return-address-with-bkey"
7678 *ValPtr = CI->getZExtValue();
7682 LLVM_DEBUG(
dbgs() <<
"Found module flag: " << IDStr <<
"(" << *ValPtr
7687 bool BTE = BTEValue == 1;
7688 bool BPPLR = BPPLRValue == 1;
7689 bool GCS = GCSValue == 1;
7690 bool SRA = SRAValue == 1;
7693 if (SRA && SRAALLValue == 1)
7694 SignTypeValue =
"all";
7697 if (SRA && SRABKeyValue == 1)
7698 SignKeyValue =
"b_key";
7700 for (
Function &
F : M.getFunctionList()) {
7701 if (
F.isDeclaration())
7708 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7709 A.isValid() &&
"none" ==
A.getValueAsString()) {
7710 F.removeFnAttr(
"sign-return-address");
7711 F.removeFnAttr(
"sign-return-address-key");
7727 if (SRAALLValue == 1)
7729 if (SRABKeyValue == 1) {
7758 if (
T->getNumOperands() < 1)
7763 if (S->getString().starts_with(
"llvm.vectorizer."))
7769 StringRef OldPrefix =
"llvm.vectorizer.";
7772 if (OldTag ==
"llvm.vectorizer.unroll")
7784 if (
T->getNumOperands() < 1)
7796 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7809 Ops.reserve(
T->getNumOperands());
7810 Ops.push_back(NewTag);
7811 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7812 Ops.push_back(
T->getOperand(
I));
7829 if (
T->isDistinct()) {
7830 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7842 Ops.reserve(
T->getNumOperands());
7853 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7854 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7855 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7858 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7860 auto I =
DL.find(
"-n64-");
7862 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7867 std::string Res =
DL.str();
7870 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7871 Res.append(Res.empty() ?
"G1" :
"-G1");
7879 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7880 Res.append(
"-ni:7:8:9");
7882 if (
DL.ends_with(
"ni:7"))
7884 if (
DL.ends_with(
"ni:7:8"))
7889 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7890 Res.append(
"-p7:160:256:256:32");
7891 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7892 Res.append(
"-p8:128:128:128:48");
7893 constexpr StringRef OldP8(
"-p8:128:128-");
7894 if (
DL.contains(OldP8))
7895 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7896 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7897 Res.append(
"-p9:192:256:256:32");
7902 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7903 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7904 Res.append((
"-" + AS +
":32:32").str());
7909 if (!
DL.contains(
"m:e"))
7910 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7915 if (
T.isSystemZ() && !
DL.empty()) {
7917 if (!
DL.contains(
"-S64"))
7918 return "E-S64" +
DL.drop_front(1).str();
7922 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7925 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7926 if (!
DL.contains(AddrSpaces)) {
7928 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7929 if (R.match(Res, &
Groups))
7935 if (
T.isAArch64()) {
7937 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7938 Res.append(
"-Fn32");
7939 AddPtr32Ptr64AddrSpaces();
7943 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7947 std::string I64 =
"-i64:64";
7948 std::string I128 =
"-i128:128";
7950 size_t Pos = Res.find(I64);
7951 if (Pos !=
size_t(-1))
7952 Res.insert(Pos + I64.size(), I128);
7956 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7957 size_t Pos = Res.find(
"-S128");
7960 Res.insert(Pos,
"-f64:32:64");
7965 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
7967 size_t Pos = Res.
find(p3232);
7969 Res.insert(Pos + p3232.
size(),
"-Fi8");
7975 AddPtr32Ptr64AddrSpaces();
7983 if (!
T.isOSIAMCU()) {
7984 std::string I128 =
"-i128:128";
7987 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7988 if (R.match(Res, &
Groups))
7996 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7998 auto I =
Ref.find(
"-f80:32-");
8000 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
8008 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
8011 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
8012 B.removeAttribute(
"no-frame-pointer-elim");
8014 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
8016 if (FramePointer !=
"all")
8017 FramePointer =
"non-leaf";
8018 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
8020 if (!FramePointer.
empty())
8021 B.addAttribute(
"frame-pointer", FramePointer);
8023 A =
B.getAttribute(
"null-pointer-is-valid");
8026 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
8027 B.removeAttribute(
"null-pointer-is-valid");
8028 if (NullPointerIsValid)
8029 B.addAttribute(Attribute::NullPointerIsValid);
8032 A =
B.getAttribute(
"uniform-work-group-size");
8036 bool IsTrue = Val ==
"true";
8037 B.removeAttribute(
"uniform-work-group-size");
8039 B.addAttribute(
"uniform-work-group-size");
8050 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 Value * upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI, StringRef Name, const Intrinsic::ID(&IIDs)[2][2])
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)
static constexpr Intrinsic::ID NVVMFAddIIDs[2][2]
#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 std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFPArithUpgrade(StringRef Name, const Intrinsic::ID(&IIDs)[2][2])
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 bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static void ConvertModuleFlag(Module &M, Module::ModFlagBehavior Behavior, StringRef Key, uint32_t Val)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static 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 constexpr Intrinsic::ID NVVMFMulIIDs[2][2]
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',...
Module * getParent()
Get the module that this global value is contained inside of...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
op_iterator op_end() const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
op_iterator op_begin() const
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
bool consume_front(char Prefix)
Returns true if this StringRef has the given prefix and removes that prefix.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
StringSwitch & StartsWith(StringLiteral S, T Value)
StringSwitch & Cases(std::initializer_list< StringLiteral > CaseStrings, T Value)
Class to represent struct types.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
unsigned getNumElements() const
Random access to the elements.
Type * getElementType(unsigned N) const
The TimeTraceScope is a helper class to call the begin and end functions of the time trace profiler.
Triple - Helper class for working with autoconf configuration names.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)
bool isVectorTy() const
True if this is an instance of VectorType.
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
LLVM_ABI unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)
LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isPtrOrPtrVectorTy() const
Return true if this is a pointer type or a vector of pointer types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isFPOrFPVectorTy() const
Return true if this is a FP type or a vector of FP.
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
static LLVM_ABI Type * getBFloatTy(LLVMContext &C)
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
A Use represents the edge between a Value definition and its users.
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
LLVM_ABI void print(raw_ostream &O, bool IsForDebug=false) const
Implement operator<< on Value.
LLVM_ABI void 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_ATTRIBUTE_VISIBILITY_DEFAULT AnalysisKey InnerAnalysisManagerProxy< AnalysisManagerT, IRUnitT, ExtraArgTs... >::Key
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
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...
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.