35#include "llvm/IR/IntrinsicsAArch64.h"
36#include "llvm/IR/IntrinsicsAMDGPU.h"
37#include "llvm/IR/IntrinsicsARM.h"
38#include "llvm/IR/IntrinsicsNVPTX.h"
39#include "llvm/IR/IntrinsicsRISCV.h"
40#include "llvm/IR/IntrinsicsWebAssembly.h"
41#include "llvm/IR/IntrinsicsX86.h"
66 cl::desc(
"Disable autoupgrade of debug info"));
85 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
100 Type *LastArgType =
F->getFunctionType()->getParamType(
101 F->getFunctionType()->getNumParams() - 1);
116 if (
F->getReturnType()->isVectorTy())
129 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
130 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
147 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
148 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
162 if (
F->getReturnType()->getScalarType()->isBFloatTy())
172 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
186 if (Name.consume_front(
"avx."))
187 return (Name.starts_with(
"blend.p") ||
188 Name ==
"cvt.ps2.pd.256" ||
189 Name ==
"cvtdq2.pd.256" ||
190 Name ==
"cvtdq2.ps.256" ||
191 Name.starts_with(
"movnt.") ||
192 Name.starts_with(
"sqrt.p") ||
193 Name.starts_with(
"storeu.") ||
194 Name.starts_with(
"vbroadcast.s") ||
195 Name.starts_with(
"vbroadcastf128") ||
196 Name.starts_with(
"vextractf128.") ||
197 Name.starts_with(
"vinsertf128.") ||
198 Name.starts_with(
"vperm2f128.") ||
199 Name.starts_with(
"vpermil."));
201 if (Name.consume_front(
"avx2."))
202 return (Name ==
"movntdqa" ||
203 Name.starts_with(
"pabs.") ||
204 Name.starts_with(
"padds.") ||
205 Name.starts_with(
"paddus.") ||
206 Name.starts_with(
"pblendd.") ||
208 Name.starts_with(
"pbroadcast") ||
209 Name.starts_with(
"pcmpeq.") ||
210 Name.starts_with(
"pcmpgt.") ||
211 Name.starts_with(
"pmax") ||
212 Name.starts_with(
"pmin") ||
213 Name.starts_with(
"pmovsx") ||
214 Name.starts_with(
"pmovzx") ||
216 Name ==
"pmulu.dq" ||
217 Name.starts_with(
"psll.dq") ||
218 Name.starts_with(
"psrl.dq") ||
219 Name.starts_with(
"psubs.") ||
220 Name.starts_with(
"psubus.") ||
221 Name.starts_with(
"vbroadcast") ||
222 Name ==
"vbroadcasti128" ||
223 Name ==
"vextracti128" ||
224 Name ==
"vinserti128" ||
225 Name ==
"vperm2i128");
227 if (Name.consume_front(
"avx512.")) {
228 if (Name.consume_front(
"mask."))
230 return (Name.starts_with(
"add.p") ||
231 Name.starts_with(
"and.") ||
232 Name.starts_with(
"andn.") ||
233 Name.starts_with(
"broadcast.s") ||
234 Name.starts_with(
"broadcastf32x4.") ||
235 Name.starts_with(
"broadcastf32x8.") ||
236 Name.starts_with(
"broadcastf64x2.") ||
237 Name.starts_with(
"broadcastf64x4.") ||
238 Name.starts_with(
"broadcasti32x4.") ||
239 Name.starts_with(
"broadcasti32x8.") ||
240 Name.starts_with(
"broadcasti64x2.") ||
241 Name.starts_with(
"broadcasti64x4.") ||
242 Name.starts_with(
"cmp.b") ||
243 Name.starts_with(
"cmp.d") ||
244 Name.starts_with(
"cmp.q") ||
245 Name.starts_with(
"cmp.w") ||
246 Name.starts_with(
"compress.b") ||
247 Name.starts_with(
"compress.d") ||
248 Name.starts_with(
"compress.p") ||
249 Name.starts_with(
"compress.q") ||
250 Name.starts_with(
"compress.store.") ||
251 Name.starts_with(
"compress.w") ||
252 Name.starts_with(
"conflict.") ||
253 Name.starts_with(
"cvtdq2pd.") ||
254 Name.starts_with(
"cvtdq2ps.") ||
255 Name ==
"cvtpd2dq.256" ||
256 Name ==
"cvtpd2ps.256" ||
257 Name ==
"cvtps2pd.128" ||
258 Name ==
"cvtps2pd.256" ||
259 Name.starts_with(
"cvtqq2pd.") ||
260 Name ==
"cvtqq2ps.256" ||
261 Name ==
"cvtqq2ps.512" ||
262 Name ==
"cvttpd2dq.256" ||
263 Name ==
"cvttps2dq.128" ||
264 Name ==
"cvttps2dq.256" ||
265 Name.starts_with(
"cvtudq2pd.") ||
266 Name.starts_with(
"cvtudq2ps.") ||
267 Name.starts_with(
"cvtuqq2pd.") ||
268 Name ==
"cvtuqq2ps.256" ||
269 Name ==
"cvtuqq2ps.512" ||
270 Name.starts_with(
"dbpsadbw.") ||
271 Name.starts_with(
"div.p") ||
272 Name.starts_with(
"expand.b") ||
273 Name.starts_with(
"expand.d") ||
274 Name.starts_with(
"expand.load.") ||
275 Name.starts_with(
"expand.p") ||
276 Name.starts_with(
"expand.q") ||
277 Name.starts_with(
"expand.w") ||
278 Name.starts_with(
"fpclass.p") ||
279 Name.starts_with(
"insert") ||
280 Name.starts_with(
"load.") ||
281 Name.starts_with(
"loadu.") ||
282 Name.starts_with(
"lzcnt.") ||
283 Name.starts_with(
"max.p") ||
284 Name.starts_with(
"min.p") ||
285 Name.starts_with(
"movddup") ||
286 Name.starts_with(
"move.s") ||
287 Name.starts_with(
"movshdup") ||
288 Name.starts_with(
"movsldup") ||
289 Name.starts_with(
"mul.p") ||
290 Name.starts_with(
"or.") ||
291 Name.starts_with(
"pabs.") ||
292 Name.starts_with(
"packssdw.") ||
293 Name.starts_with(
"packsswb.") ||
294 Name.starts_with(
"packusdw.") ||
295 Name.starts_with(
"packuswb.") ||
296 Name.starts_with(
"padd.") ||
297 Name.starts_with(
"padds.") ||
298 Name.starts_with(
"paddus.") ||
299 Name.starts_with(
"palignr.") ||
300 Name.starts_with(
"pand.") ||
301 Name.starts_with(
"pandn.") ||
302 Name.starts_with(
"pavg") ||
303 Name.starts_with(
"pbroadcast") ||
304 Name.starts_with(
"pcmpeq.") ||
305 Name.starts_with(
"pcmpgt.") ||
306 Name.starts_with(
"perm.df.") ||
307 Name.starts_with(
"perm.di.") ||
308 Name.starts_with(
"permvar.") ||
309 Name.starts_with(
"pmaddubs.w.") ||
310 Name.starts_with(
"pmaddw.d.") ||
311 Name.starts_with(
"pmax") ||
312 Name.starts_with(
"pmin") ||
313 Name ==
"pmov.qd.256" ||
314 Name ==
"pmov.qd.512" ||
315 Name ==
"pmov.wb.256" ||
316 Name ==
"pmov.wb.512" ||
317 Name.starts_with(
"pmovsx") ||
318 Name.starts_with(
"pmovzx") ||
319 Name.starts_with(
"pmul.dq.") ||
320 Name.starts_with(
"pmul.hr.sw.") ||
321 Name.starts_with(
"pmulh.w.") ||
322 Name.starts_with(
"pmulhu.w.") ||
323 Name.starts_with(
"pmull.") ||
324 Name.starts_with(
"pmultishift.qb.") ||
325 Name.starts_with(
"pmulu.dq.") ||
326 Name.starts_with(
"por.") ||
327 Name.starts_with(
"prol.") ||
328 Name.starts_with(
"prolv.") ||
329 Name.starts_with(
"pror.") ||
330 Name.starts_with(
"prorv.") ||
331 Name.starts_with(
"pshuf.b.") ||
332 Name.starts_with(
"pshuf.d.") ||
333 Name.starts_with(
"pshufh.w.") ||
334 Name.starts_with(
"pshufl.w.") ||
335 Name.starts_with(
"psll.d") ||
336 Name.starts_with(
"psll.q") ||
337 Name.starts_with(
"psll.w") ||
338 Name.starts_with(
"pslli") ||
339 Name.starts_with(
"psllv") ||
340 Name.starts_with(
"psra.d") ||
341 Name.starts_with(
"psra.q") ||
342 Name.starts_with(
"psra.w") ||
343 Name.starts_with(
"psrai") ||
344 Name.starts_with(
"psrav") ||
345 Name.starts_with(
"psrl.d") ||
346 Name.starts_with(
"psrl.q") ||
347 Name.starts_with(
"psrl.w") ||
348 Name.starts_with(
"psrli") ||
349 Name.starts_with(
"psrlv") ||
350 Name.starts_with(
"psub.") ||
351 Name.starts_with(
"psubs.") ||
352 Name.starts_with(
"psubus.") ||
353 Name.starts_with(
"pternlog.") ||
354 Name.starts_with(
"punpckh") ||
355 Name.starts_with(
"punpckl") ||
356 Name.starts_with(
"pxor.") ||
357 Name.starts_with(
"shuf.f") ||
358 Name.starts_with(
"shuf.i") ||
359 Name.starts_with(
"shuf.p") ||
360 Name.starts_with(
"sqrt.p") ||
361 Name.starts_with(
"store.b.") ||
362 Name.starts_with(
"store.d.") ||
363 Name.starts_with(
"store.p") ||
364 Name.starts_with(
"store.q.") ||
365 Name.starts_with(
"store.w.") ||
366 Name ==
"store.ss" ||
367 Name.starts_with(
"storeu.") ||
368 Name.starts_with(
"sub.p") ||
369 Name.starts_with(
"ucmp.") ||
370 Name.starts_with(
"unpckh.") ||
371 Name.starts_with(
"unpckl.") ||
372 Name.starts_with(
"valign.") ||
373 Name ==
"vcvtph2ps.128" ||
374 Name ==
"vcvtph2ps.256" ||
375 Name.starts_with(
"vextract") ||
376 Name.starts_with(
"vfmadd.") ||
377 Name.starts_with(
"vfmaddsub.") ||
378 Name.starts_with(
"vfnmadd.") ||
379 Name.starts_with(
"vfnmsub.") ||
380 Name.starts_with(
"vpdpbusd.") ||
381 Name.starts_with(
"vpdpbusds.") ||
382 Name.starts_with(
"vpdpwssd.") ||
383 Name.starts_with(
"vpdpwssds.") ||
384 Name.starts_with(
"vpermi2var.") ||
385 Name.starts_with(
"vpermil.p") ||
386 Name.starts_with(
"vpermilvar.") ||
387 Name.starts_with(
"vpermt2var.") ||
388 Name.starts_with(
"vpmadd52") ||
389 Name.starts_with(
"vpshld.") ||
390 Name.starts_with(
"vpshldv.") ||
391 Name.starts_with(
"vpshrd.") ||
392 Name.starts_with(
"vpshrdv.") ||
393 Name.starts_with(
"vpshufbitqmb.") ||
394 Name.starts_with(
"xor."));
396 if (Name.consume_front(
"mask3."))
398 return (Name.starts_with(
"vfmadd.") ||
399 Name.starts_with(
"vfmaddsub.") ||
400 Name.starts_with(
"vfmsub.") ||
401 Name.starts_with(
"vfmsubadd.") ||
402 Name.starts_with(
"vfnmsub."));
404 if (Name.consume_front(
"maskz."))
406 return (Name.starts_with(
"pternlog.") ||
407 Name.starts_with(
"vfmadd.") ||
408 Name.starts_with(
"vfmaddsub.") ||
409 Name.starts_with(
"vpdpbusd.") ||
410 Name.starts_with(
"vpdpbusds.") ||
411 Name.starts_with(
"vpdpwssd.") ||
412 Name.starts_with(
"vpdpwssds.") ||
413 Name.starts_with(
"vpermt2var.") ||
414 Name.starts_with(
"vpmadd52") ||
415 Name.starts_with(
"vpshldv.") ||
416 Name.starts_with(
"vpshrdv."));
419 return (Name ==
"movntdqa" ||
420 Name ==
"pmul.dq.512" ||
421 Name ==
"pmulu.dq.512" ||
422 Name.starts_with(
"broadcastm") ||
423 Name.starts_with(
"cmp.p") ||
424 Name.starts_with(
"cvtb2mask.") ||
425 Name.starts_with(
"cvtd2mask.") ||
426 Name.starts_with(
"cvtmask2") ||
427 Name.starts_with(
"cvtq2mask.") ||
428 Name ==
"cvtusi2sd" ||
429 Name.starts_with(
"cvtw2mask.") ||
434 Name ==
"kortestc.w" ||
435 Name ==
"kortestz.w" ||
436 Name.starts_with(
"kunpck") ||
439 Name.starts_with(
"padds.") ||
440 Name.starts_with(
"pbroadcast") ||
441 Name.starts_with(
"prol") ||
442 Name.starts_with(
"pror") ||
443 Name.starts_with(
"psll.dq") ||
444 Name.starts_with(
"psrl.dq") ||
445 Name.starts_with(
"psubs.") ||
446 Name.starts_with(
"ptestm") ||
447 Name.starts_with(
"ptestnm") ||
448 Name.starts_with(
"storent.") ||
449 Name.starts_with(
"vbroadcast.s") ||
450 Name.starts_with(
"vpshld.") ||
451 Name.starts_with(
"vpshrd."));
454 if (Name.consume_front(
"fma."))
455 return (Name.starts_with(
"vfmadd.") ||
456 Name.starts_with(
"vfmsub.") ||
457 Name.starts_with(
"vfmsubadd.") ||
458 Name.starts_with(
"vfnmadd.") ||
459 Name.starts_with(
"vfnmsub."));
461 if (Name.consume_front(
"fma4."))
462 return Name.starts_with(
"vfmadd.s");
464 if (Name.consume_front(
"sse."))
465 return (Name ==
"add.ss" ||
466 Name ==
"cvtsi2ss" ||
467 Name ==
"cvtsi642ss" ||
470 Name.starts_with(
"sqrt.p") ||
472 Name.starts_with(
"storeu.") ||
475 if (Name.consume_front(
"sse2."))
476 return (Name ==
"add.sd" ||
477 Name ==
"cvtdq2pd" ||
478 Name ==
"cvtdq2ps" ||
479 Name ==
"cvtps2pd" ||
480 Name ==
"cvtsi2sd" ||
481 Name ==
"cvtsi642sd" ||
482 Name ==
"cvtss2sd" ||
485 Name.starts_with(
"padds.") ||
486 Name.starts_with(
"paddus.") ||
487 Name.starts_with(
"pcmpeq.") ||
488 Name.starts_with(
"pcmpgt.") ||
493 Name ==
"pmulu.dq" ||
494 Name.starts_with(
"pshuf") ||
495 Name.starts_with(
"psll.dq") ||
496 Name.starts_with(
"psrl.dq") ||
497 Name.starts_with(
"psubs.") ||
498 Name.starts_with(
"psubus.") ||
499 Name.starts_with(
"sqrt.p") ||
501 Name ==
"storel.dq" ||
502 Name.starts_with(
"storeu.") ||
505 if (Name.consume_front(
"sse41."))
506 return (Name.starts_with(
"blendp") ||
507 Name ==
"movntdqa" ||
517 Name.starts_with(
"pmovsx") ||
518 Name.starts_with(
"pmovzx") ||
521 if (Name.consume_front(
"sse42."))
522 return Name ==
"crc32.64.8";
524 if (Name.consume_front(
"sse4a."))
525 return Name.starts_with(
"movnt.");
527 if (Name.consume_front(
"ssse3."))
528 return (Name ==
"pabs.b.128" ||
529 Name ==
"pabs.d.128" ||
530 Name ==
"pabs.w.128");
532 if (Name.consume_front(
"xop."))
533 return (Name ==
"vpcmov" ||
534 Name ==
"vpcmov.256" ||
535 Name.starts_with(
"vpcom") ||
536 Name.starts_with(
"vprot"));
538 if (Name.consume_front(
"bmi."))
539 return (Name.starts_with(
"pdep.") ||
540 Name.starts_with(
"pext."));
542 return (Name ==
"addcarry.u32" ||
543 Name ==
"addcarry.u64" ||
544 Name ==
"addcarryx.u32" ||
545 Name ==
"addcarryx.u64" ||
546 Name ==
"subborrow.u32" ||
547 Name ==
"subborrow.u64" ||
548 Name.starts_with(
"vcvtph2ps."));
554 if (!Name.consume_front(
"x86."))
562 if (Name ==
"rdtscp") {
564 if (
F->getFunctionType()->getNumParams() == 0)
569 Intrinsic::x86_rdtscp);
576 if (Name.consume_front(
"sse41.ptest")) {
578 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
579 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
580 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
593 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
594 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
595 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
596 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
597 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
598 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
603 if (Name.consume_front(
"avx512.")) {
604 if (Name.consume_front(
"mask.cmp.")) {
607 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
608 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
609 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
610 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
611 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
612 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
616 }
else if (Name.starts_with(
"vpdpbusd.") ||
617 Name.starts_with(
"vpdpbusds.")) {
620 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
621 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
622 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
623 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
624 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
625 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
629 }
else if (Name.starts_with(
"vpdpwssd.") ||
630 Name.starts_with(
"vpdpwssds.")) {
633 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
634 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
635 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
636 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
637 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
638 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
646 if (Name.consume_front(
"avx2.")) {
647 if (Name.consume_front(
"vpdpb")) {
650 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
651 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
652 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
653 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
654 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
655 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
656 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
657 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
658 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
659 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
660 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
661 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
665 }
else if (Name.consume_front(
"vpdpw")) {
668 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
669 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
670 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
671 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
672 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
673 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
674 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
675 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
676 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
677 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
678 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
679 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
687 if (Name.consume_front(
"avx10.")) {
688 if (Name.consume_front(
"vpdpb")) {
691 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
692 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
693 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
694 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
695 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
696 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
700 }
else if (Name.consume_front(
"vpdpw")) {
702 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
703 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
704 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
705 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
706 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
707 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
715 if (Name.consume_front(
"avx512bf16.")) {
718 .
Case(
"cvtne2ps2bf16.128",
719 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
720 .
Case(
"cvtne2ps2bf16.256",
721 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
722 .
Case(
"cvtne2ps2bf16.512",
723 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
724 .
Case(
"mask.cvtneps2bf16.128",
725 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
726 .
Case(
"cvtneps2bf16.256",
727 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
728 .
Case(
"cvtneps2bf16.512",
729 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
736 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
737 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
738 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
745 if (Name.consume_front(
"xop.")) {
747 if (Name.starts_with(
"vpermil2")) {
750 auto Idx =
F->getFunctionType()->getParamType(2);
751 if (Idx->isFPOrFPVectorTy()) {
752 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
753 unsigned EltSize = Idx->getScalarSizeInBits();
754 if (EltSize == 64 && IdxSize == 128)
755 ID = Intrinsic::x86_xop_vpermil2pd;
756 else if (EltSize == 32 && IdxSize == 128)
757 ID = Intrinsic::x86_xop_vpermil2ps;
758 else if (EltSize == 64 && IdxSize == 256)
759 ID = Intrinsic::x86_xop_vpermil2pd_256;
761 ID = Intrinsic::x86_xop_vpermil2ps_256;
763 }
else if (
F->arg_size() == 2)
766 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
767 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
778 if (Name ==
"seh.recoverfp") {
780 Intrinsic::eh_recoverfp);
792 if (Name.starts_with(
"rbit")) {
795 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
799 if (Name ==
"thread.pointer") {
802 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
806 bool Neon = Name.consume_front(
"neon.");
811 if (Name.consume_front(
"bfdot.")) {
815 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
820 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
821 assert((OperandWidth == 64 || OperandWidth == 128) &&
822 "Unexpected operand width");
824 std::array<Type *, 2> Tys{
835 if (Name.consume_front(
"bfm")) {
837 if (Name.consume_back(
".v4f32.v16i8")) {
883 F->arg_begin()->getType());
887 if (Name.consume_front(
"vst")) {
889 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
893 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
894 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
897 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
898 Intrinsic::arm_neon_vst4lane};
900 auto fArgs =
F->getFunctionType()->params();
901 Type *Tys[] = {fArgs[0], fArgs[1]};
904 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
907 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
916 if (Name.consume_front(
"mve.")) {
918 if (Name ==
"vctp64") {
928 if (Name.starts_with(
"vrintn.v")) {
930 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
935 if (Name.consume_back(
".v4i1")) {
937 if (Name.consume_back(
".predicated.v2i64.v4i32"))
939 return Name ==
"mull.int" || Name ==
"vqdmull";
941 if (Name.consume_back(
".v2i64")) {
943 bool IsGather = Name.consume_front(
"vldr.gather.");
944 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
945 if (Name.consume_front(
"base.")) {
947 Name.consume_front(
"wb.");
950 return Name ==
"predicated.v2i64";
953 if (Name.consume_front(
"offset.predicated."))
954 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
955 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
968 if (Name.consume_front(
"cde.vcx")) {
970 if (Name.consume_back(
".predicated.v2i64.v4i1"))
972 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
973 Name ==
"3q" || Name ==
"3qa";
987 F->arg_begin()->getType());
991 if (Name.starts_with(
"addp")) {
993 if (
F->arg_size() != 2)
996 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
998 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1004 if (Name.starts_with(
"bfcvt")) {
1010 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1017 if (Name.consume_front(
"sve.")) {
1019 if (Name.consume_front(
"bf")) {
1020 if (Name ==
"mmla") {
1021 Type *Tys[] = {
F->getReturnType(),
1022 std::next(
F->arg_begin())->getType()};
1024 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1027 if (Name.consume_back(
".lane")) {
1031 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1032 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1033 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1045 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1050 if (Name.consume_front(
"convert.from.svbool")) {
1053 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1056 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1061 if (Name.consume_front(
"convert.to.svbool")) {
1064 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1067 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1072 if (Name.consume_front(
"addqv")) {
1074 if (!
F->getReturnType()->isFPOrFPVectorTy())
1077 auto Args =
F->getFunctionType()->params();
1078 Type *Tys[] = {
F->getReturnType(), Args[1]};
1080 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1084 if (Name.consume_front(
"ld")) {
1086 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1087 if (LdRegex.
match(Name)) {
1093 "Expected 2 arguments for ld* intrinsic.");
1094 Type *PtrTy =
F->getArg(1)->getType();
1097 Intrinsic::aarch64_sve_ld2_sret,
1098 Intrinsic::aarch64_sve_ld3_sret,
1099 Intrinsic::aarch64_sve_ld4_sret,
1102 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1108 if (Name.consume_front(
"tuple.")) {
1110 if (Name.starts_with(
"get")) {
1112 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1114 F->getParent(), Intrinsic::vector_extract, Tys);
1118 if (Name.starts_with(
"set")) {
1120 auto Args =
F->getFunctionType()->params();
1121 Type *Tys[] = {Args[0], Args[2], Args[1]};
1123 F->getParent(), Intrinsic::vector_insert, Tys);
1127 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1128 if (CreateTupleRegex.
match(Name)) {
1130 auto Args =
F->getFunctionType()->params();
1131 Type *Tys[] = {
F->getReturnType(), Args[1]};
1133 F->getParent(), Intrinsic::vector_insert, Tys);
1139 if (Name.starts_with(
"rev.nxv")) {
1142 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1148 if (Name.consume_front(
"sme.")) {
1150 if (Name.consume_front(
"ftmopa.")) {
1155 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1156 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1173 if (Name.consume_front(
"cp.async.bulk.tensor.g2s.")) {
1177 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d)
1179 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d)
1181 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d)
1182 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d)
1183 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d)
1184 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d)
1185 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d)
1186 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d)
1195 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1209 size_t FlagStartIndex =
F->getFunctionType()->getNumParams() - 3;
1210 Type *ArgType =
F->getFunctionType()->getParamType(FlagStartIndex);
1235 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1238 auto [RedOpName, ShapeName] = Name.split(
'.');
1243 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1244 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1245 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1246 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1247 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1248 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1249 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1250 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1256 if (Name.consume_front(
"mapa.shared.cluster"))
1257 if (
F->getReturnType()->getPointerAddressSpace() ==
1259 return Intrinsic::nvvm_mapa_shared_cluster;
1261 if (Name.consume_front(
"cp.async.bulk.")) {
1264 .
Case(
"global.to.shared.cluster",
1265 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1266 .
Case(
"shared.cta.to.cluster",
1267 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1271 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1281 if (!Name.consume_front(
"tcgen05.commit."))
1284 if (Name.consume_front(
"shared."))
1286 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1287 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1290 if (Name.consume_front(
"mc.shared.")) {
1292 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1296 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1297 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1306 if (
F->arg_size() != 2)
1309 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1310 Name.consume_front(
"tcgen05.alloc."))
1312 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1313 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1316 if (Name.consume_front(
"tcgen05.dealloc."))
1318 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1319 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1326 if (Name.consume_front(
"fma.rn."))
1328 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1329 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1330 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1331 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1334 if (Name.consume_front(
"fmax."))
1336 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1337 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1338 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1339 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1340 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1341 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1342 .
Case(
"ftz.nan.xorsign.abs.bf16",
1343 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1344 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1345 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1346 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1347 .
Case(
"ftz.xorsign.abs.bf16x2",
1348 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1349 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1350 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1351 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1352 .
Case(
"nan.xorsign.abs.bf16x2",
1353 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1354 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1355 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1358 if (Name.consume_front(
"fmin."))
1360 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1361 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1362 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1363 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1364 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1365 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1366 .
Case(
"ftz.nan.xorsign.abs.bf16",
1367 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1368 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1369 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1370 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1371 .
Case(
"ftz.xorsign.abs.bf16x2",
1372 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1373 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1374 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1375 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1376 .
Case(
"nan.xorsign.abs.bf16x2",
1377 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1378 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1379 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1382 if (Name.consume_front(
"neg."))
1384 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1385 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1393 if (!Name.consume_front(
"tcgen05.mma."))
1397 if (Name.starts_with(
"ws"))
1400 return F->getIntrinsicID();
1404 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1405 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1406 Name.consume_front(
"param");
1410 if (!Name.consume_front(
"vp."))
1439 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1440 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1447 if (!Name.consume_front(
"vp."))
1467 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1468 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1473 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1485 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1496 if (Name.starts_with(
"to.fp16")) {
1500 FuncTy->getReturnType());
1503 if (Name.starts_with(
"from.fp16")) {
1507 FuncTy->getReturnType());
1519 if (Defaults.empty())
1531 if (
F->arg_size() >= FullDecl->
arg_size())
1536 if (
F->arg_size() < FirstDefault)
1544 bool CanUpgradeDebugIntrinsicsToRecords) {
1545 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1550 if (!Name.consume_front(
"llvm.") || Name.empty())
1556 bool IsArm = Name.consume_front(
"arm.");
1557 if (IsArm || Name.consume_front(
"aarch64.")) {
1563 if (Name.consume_front(
"amdgcn.")) {
1564 if (Name ==
"alignbit") {
1567 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1571 if (Name.consume_front(
"atomic.")) {
1572 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1573 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1582 switch (
F->getIntrinsicID()) {
1586 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1587 if (
F->arg_size() == 7) {
1592 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1593 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1594 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1595 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1596 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1597 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1598 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1599 if (
F->arg_size() == 8) {
1606 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1607 Name.consume_front(
"flat.atomic.")) {
1608 if (Name.starts_with(
"fadd") ||
1610 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1611 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1619 if (Name.starts_with(
"ldexp.")) {
1622 F->getParent(), Intrinsic::ldexp,
1623 {F->getReturnType(), F->getArg(1)->getType()});
1632 if (
F->arg_size() == 1) {
1633 if (Name.consume_front(
"convert.")) {
1647 F->arg_begin()->getType());
1653 if (Name ==
"coro.end" &&
1654 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1655 CoroEndID = Intrinsic::coro_end;
1656 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1657 CoroEndID = Intrinsic::coro_end_async;
1668 if (Name.consume_front(
"dbg.")) {
1670 if (CanUpgradeDebugIntrinsicsToRecords) {
1671 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1672 Name ==
"declare" || Name ==
"label") {
1681 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1684 Intrinsic::dbg_value);
1691 if (Name.consume_front(
"experimental.vector.")) {
1697 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1698 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1699 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1700 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1701 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1703 Intrinsic::vector_partial_reduce_add)
1706 const auto *FT =
F->getFunctionType();
1708 if (ID == Intrinsic::vector_extract ||
1709 ID == Intrinsic::vector_interleave2)
1712 if (ID != Intrinsic::vector_interleave2)
1714 if (ID == Intrinsic::vector_insert ||
1715 ID == Intrinsic::vector_partial_reduce_add)
1723 if (Name.consume_front(
"reduce.")) {
1725 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1726 if (R.match(Name, &
Groups))
1728 .
Case(
"add", Intrinsic::vector_reduce_add)
1729 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1730 .
Case(
"and", Intrinsic::vector_reduce_and)
1731 .
Case(
"or", Intrinsic::vector_reduce_or)
1732 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1733 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1734 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1735 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1736 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1737 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1738 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1743 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1748 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1749 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1754 auto Args =
F->getFunctionType()->params();
1756 {Args[V2 ? 1 : 0]});
1762 if (Name.consume_front(
"splice"))
1766 if (Name.consume_front(
"experimental.stepvector.")) {
1770 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1775 if (Name.starts_with(
"flt.rounds")) {
1778 Intrinsic::get_rounding);
1783 if (Name.starts_with(
"invariant.group.barrier")) {
1785 auto Args =
F->getFunctionType()->params();
1786 Type* ObjectPtr[1] = {Args[0]};
1789 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1794 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1795 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1796 if (IsLifetimeStart || IsLifetimeEnd) {
1797 if (
F->arg_size() == 2) {
1798 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1799 : Intrinsic::lifetime_end;
1804 F->getArg(1)->getType());
1806 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1826 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1827 .StartsWith(
"memmove.", Intrinsic::memmove)
1829 if (
F->arg_size() == 5) {
1833 F->getFunctionType()->params().slice(0, 3);
1839 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1842 const auto *FT =
F->getFunctionType();
1843 Type *ParamTypes[2] = {
1844 FT->getParamType(0),
1848 Intrinsic::memset, ParamTypes);
1854 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1855 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1856 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1857 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1859 if (MaskedID &&
F->arg_size() == 4) {
1861 if (MaskedID == Intrinsic::masked_load ||
1862 MaskedID == Intrinsic::masked_gather) {
1864 F->getParent(), MaskedID,
1865 {F->getReturnType(), F->getArg(0)->getType()});
1869 F->getParent(), MaskedID,
1870 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1876 if (Name.consume_front(
"nvvm.")) {
1878 if (
F->arg_size() == 1) {
1881 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1882 .Case(
"clz.i", Intrinsic::ctlz)
1883 .
Case(
"popc.i", Intrinsic::ctpop)
1887 {F->getReturnType()});
1890 }
else if (
F->arg_size() == 2) {
1893 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1894 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1895 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1896 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1900 {F->getReturnType()});
1906 if (!
F->getReturnType()->getScalarType()->isBFloatTy()) {
1936 F->getParent(), IID,
F->getReturnType(),
1937 F->getFunctionType()->params());
1948 {F->getArg(0)->getType()});
1973 bool Expand =
false;
1974 if (Name.consume_front(
"abs."))
1977 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
1978 else if (Name.consume_front(
"fabs."))
1980 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
1981 else if (Name.consume_front(
"ex2.approx."))
1984 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
1985 else if (Name.consume_front(
"atomic.load."))
1994 else if (Name.consume_front(
"atomic."))
2009 else if (Name.consume_front(
"bitcast."))
2012 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2013 else if (Name.consume_front(
"rotate."))
2015 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2016 else if (Name.consume_front(
"ptr.gen.to."))
2019 else if (Name.consume_front(
"ptr."))
2022 else if (Name.consume_front(
"ldg.global."))
2024 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2025 Name.starts_with(
"p."));
2028 .
Case(
"barrier0",
true)
2029 .
Case(
"barrier.n",
true)
2030 .
Case(
"barrier.sync.cnt",
true)
2031 .
Case(
"barrier.sync",
true)
2032 .
Case(
"barrier",
true)
2033 .
Case(
"bar.sync",
true)
2034 .
Case(
"barrier0.popc",
true)
2035 .
Case(
"barrier0.and",
true)
2036 .
Case(
"barrier0.or",
true)
2037 .
Case(
"clz.ll",
true)
2038 .
Case(
"popc.ll",
true)
2040 .
Case(
"swap.lo.hi.b64",
true)
2041 .
Case(
"tanh.approx.f32",
true)
2053 if (Name.starts_with(
"objectsize.")) {
2054 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2055 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2058 Intrinsic::objectsize, Tys);
2065 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2068 F->getParent(), Intrinsic::ptr_annotation,
2069 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2075 if (Name.consume_front(
"riscv.")) {
2078 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2079 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2080 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2081 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2084 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2097 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2098 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2107 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2108 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2109 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2110 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2115 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2124 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2126 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2135 if (Name ==
"stackprotectorcheck") {
2142 if (Name ==
"thread.pointer") {
2144 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2150 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2153 F->getParent(), Intrinsic::var_annotation,
2154 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2157 if (Name.consume_front(
"vector.splice")) {
2158 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2168 if (Name.consume_front(
"wasm.")) {
2171 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2172 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2173 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2178 F->getReturnType());
2182 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2184 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2186 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2205 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2215 std::string
Name =
F->getName().str();
2218 Name,
F->getParent());
2229 if (Result != std::nullopt) {
2245 bool CanUpgradeDebugIntrinsicsToRecords) {
2265 GV->
getName() ==
"llvm.global_dtors")) ||
2280 unsigned N =
Init->getNumOperands();
2281 std::vector<Constant *> NewCtors(
N);
2282 for (
unsigned i = 0; i !=
N; ++i) {
2285 Ctor->getAggregateElement(1),
2299 unsigned NumElts = ResultTy->getNumElements() * 8;
2303 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2313 for (
unsigned l = 0; l != NumElts; l += 16)
2314 for (
unsigned i = 0; i != 16; ++i) {
2315 unsigned Idx = NumElts + i - Shift;
2317 Idx -= NumElts - 16;
2318 Idxs[l + i] = Idx + l;
2321 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2325 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2333 unsigned NumElts = ResultTy->getNumElements() * 8;
2337 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2347 for (
unsigned l = 0; l != NumElts; l += 16)
2348 for (
unsigned i = 0; i != 16; ++i) {
2349 unsigned Idx = i + Shift;
2351 Idx += NumElts - 16;
2352 Idxs[l + i] = Idx + l;
2355 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2359 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2367 Mask = Builder.CreateBitCast(Mask, MaskTy);
2373 for (
unsigned i = 0; i != NumElts; ++i)
2375 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2386 if (
C->isAllOnesValue())
2391 return Builder.CreateSelect(Mask, Op0, Op1);
2398 if (
C->isAllOnesValue())
2402 Mask->getType()->getIntegerBitWidth());
2403 Mask = Builder.CreateBitCast(Mask, MaskTy);
2404 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2405 return Builder.CreateSelect(Mask, Op0, Op1);
2418 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2419 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2424 ShiftVal &= (NumElts - 1);
2433 if (ShiftVal > 16) {
2441 for (
unsigned l = 0; l < NumElts; l += 16) {
2442 for (
unsigned i = 0; i != 16; ++i) {
2443 unsigned Idx = ShiftVal + i;
2444 if (!IsVALIGN && Idx >= 16)
2445 Idx += NumElts - 16;
2446 Indices[l + i] = Idx + l;
2451 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2457 bool ZeroMask,
bool IndexForm) {
2460 unsigned EltWidth = Ty->getScalarSizeInBits();
2461 bool IsFloat = Ty->isFPOrFPVectorTy();
2463 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2464 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2465 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2466 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2467 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2468 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2469 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2470 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2471 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2472 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2473 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2474 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2475 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2476 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2477 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2478 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2479 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2480 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2481 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2482 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2483 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2484 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2485 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2486 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2487 else if (VecWidth == 128 && EltWidth == 16)
2488 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2489 else if (VecWidth == 256 && EltWidth == 16)
2490 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2491 else if (VecWidth == 512 && EltWidth == 16)
2492 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2493 else if (VecWidth == 128 && EltWidth == 8)
2494 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2495 else if (VecWidth == 256 && EltWidth == 8)
2496 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2497 else if (VecWidth == 512 && EltWidth == 8)
2498 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2509 Value *V = Builder.CreateIntrinsic(IID, Args);
2521 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2532 bool IsRotateRight) {
2542 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2543 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2546 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2547 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2592 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2597 bool IsShiftRight,
bool ZeroMask) {
2611 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2612 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2615 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2616 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2631 const Align Alignment =
2633 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2638 if (
C->isAllOnesValue())
2639 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2644 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2650 const Align Alignment =
2659 if (
C->isAllOnesValue())
2660 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2665 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2671 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2672 {Op0, Builder.getInt1(
false)});
2687 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2688 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2689 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2690 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2691 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2694 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2695 LHS = Builder.CreateAnd(
LHS, Mask);
2696 RHS = Builder.CreateAnd(
RHS, Mask);
2713 if (!
C || !
C->isAllOnesValue())
2714 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2719 for (
unsigned i = 0; i != NumElts; ++i)
2721 for (
unsigned i = NumElts; i != 8; ++i)
2722 Indices[i] = NumElts + i % NumElts;
2723 Vec = Builder.CreateShuffleVector(Vec,
2727 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2731 unsigned CC,
bool Signed) {
2739 }
else if (CC == 7) {
2775 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2776 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2778 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2779 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2788 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2794 Name = Name.substr(12);
2799 if (Name.starts_with(
"max.p")) {
2800 if (VecWidth == 128 && EltWidth == 32)
2801 IID = Intrinsic::x86_sse_max_ps;
2802 else if (VecWidth == 128 && EltWidth == 64)
2803 IID = Intrinsic::x86_sse2_max_pd;
2804 else if (VecWidth == 256 && EltWidth == 32)
2805 IID = Intrinsic::x86_avx_max_ps_256;
2806 else if (VecWidth == 256 && EltWidth == 64)
2807 IID = Intrinsic::x86_avx_max_pd_256;
2810 }
else if (Name.starts_with(
"min.p")) {
2811 if (VecWidth == 128 && EltWidth == 32)
2812 IID = Intrinsic::x86_sse_min_ps;
2813 else if (VecWidth == 128 && EltWidth == 64)
2814 IID = Intrinsic::x86_sse2_min_pd;
2815 else if (VecWidth == 256 && EltWidth == 32)
2816 IID = Intrinsic::x86_avx_min_ps_256;
2817 else if (VecWidth == 256 && EltWidth == 64)
2818 IID = Intrinsic::x86_avx_min_pd_256;
2821 }
else if (Name.starts_with(
"pshuf.b.")) {
2822 if (VecWidth == 128)
2823 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2824 else if (VecWidth == 256)
2825 IID = Intrinsic::x86_avx2_pshuf_b;
2826 else if (VecWidth == 512)
2827 IID = Intrinsic::x86_avx512_pshuf_b_512;
2830 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2831 if (VecWidth == 128)
2832 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2833 else if (VecWidth == 256)
2834 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2835 else if (VecWidth == 512)
2836 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2839 }
else if (Name.starts_with(
"pmulh.w.")) {
2840 if (VecWidth == 128)
2841 IID = Intrinsic::x86_sse2_pmulh_w;
2842 else if (VecWidth == 256)
2843 IID = Intrinsic::x86_avx2_pmulh_w;
2844 else if (VecWidth == 512)
2845 IID = Intrinsic::x86_avx512_pmulh_w_512;
2848 }
else if (Name.starts_with(
"pmulhu.w.")) {
2849 if (VecWidth == 128)
2850 IID = Intrinsic::x86_sse2_pmulhu_w;
2851 else if (VecWidth == 256)
2852 IID = Intrinsic::x86_avx2_pmulhu_w;
2853 else if (VecWidth == 512)
2854 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2857 }
else if (Name.starts_with(
"pmaddw.d.")) {
2858 if (VecWidth == 128)
2859 IID = Intrinsic::x86_sse2_pmadd_wd;
2860 else if (VecWidth == 256)
2861 IID = Intrinsic::x86_avx2_pmadd_wd;
2862 else if (VecWidth == 512)
2863 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2866 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2867 if (VecWidth == 128)
2868 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2869 else if (VecWidth == 256)
2870 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2871 else if (VecWidth == 512)
2872 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2875 }
else if (Name.starts_with(
"packsswb.")) {
2876 if (VecWidth == 128)
2877 IID = Intrinsic::x86_sse2_packsswb_128;
2878 else if (VecWidth == 256)
2879 IID = Intrinsic::x86_avx2_packsswb;
2880 else if (VecWidth == 512)
2881 IID = Intrinsic::x86_avx512_packsswb_512;
2884 }
else if (Name.starts_with(
"packssdw.")) {
2885 if (VecWidth == 128)
2886 IID = Intrinsic::x86_sse2_packssdw_128;
2887 else if (VecWidth == 256)
2888 IID = Intrinsic::x86_avx2_packssdw;
2889 else if (VecWidth == 512)
2890 IID = Intrinsic::x86_avx512_packssdw_512;
2893 }
else if (Name.starts_with(
"packuswb.")) {
2894 if (VecWidth == 128)
2895 IID = Intrinsic::x86_sse2_packuswb_128;
2896 else if (VecWidth == 256)
2897 IID = Intrinsic::x86_avx2_packuswb;
2898 else if (VecWidth == 512)
2899 IID = Intrinsic::x86_avx512_packuswb_512;
2902 }
else if (Name.starts_with(
"packusdw.")) {
2903 if (VecWidth == 128)
2904 IID = Intrinsic::x86_sse41_packusdw;
2905 else if (VecWidth == 256)
2906 IID = Intrinsic::x86_avx2_packusdw;
2907 else if (VecWidth == 512)
2908 IID = Intrinsic::x86_avx512_packusdw_512;
2911 }
else if (Name.starts_with(
"vpermilvar.")) {
2912 if (VecWidth == 128 && EltWidth == 32)
2913 IID = Intrinsic::x86_avx_vpermilvar_ps;
2914 else if (VecWidth == 128 && EltWidth == 64)
2915 IID = Intrinsic::x86_avx_vpermilvar_pd;
2916 else if (VecWidth == 256 && EltWidth == 32)
2917 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2918 else if (VecWidth == 256 && EltWidth == 64)
2919 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2920 else if (VecWidth == 512 && EltWidth == 32)
2921 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2922 else if (VecWidth == 512 && EltWidth == 64)
2923 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2926 }
else if (Name ==
"cvtpd2dq.256") {
2927 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2928 }
else if (Name ==
"cvtpd2ps.256") {
2929 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2930 }
else if (Name ==
"cvttpd2dq.256") {
2931 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2932 }
else if (Name ==
"cvttps2dq.128") {
2933 IID = Intrinsic::x86_sse2_cvttps2dq;
2934 }
else if (Name ==
"cvttps2dq.256") {
2935 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2936 }
else if (Name.starts_with(
"permvar.")) {
2938 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2939 IID = Intrinsic::x86_avx2_permps;
2940 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2941 IID = Intrinsic::x86_avx2_permd;
2942 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2943 IID = Intrinsic::x86_avx512_permvar_df_256;
2944 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2945 IID = Intrinsic::x86_avx512_permvar_di_256;
2946 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2947 IID = Intrinsic::x86_avx512_permvar_sf_512;
2948 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2949 IID = Intrinsic::x86_avx512_permvar_si_512;
2950 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2951 IID = Intrinsic::x86_avx512_permvar_df_512;
2952 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2953 IID = Intrinsic::x86_avx512_permvar_di_512;
2954 else if (VecWidth == 128 && EltWidth == 16)
2955 IID = Intrinsic::x86_avx512_permvar_hi_128;
2956 else if (VecWidth == 256 && EltWidth == 16)
2957 IID = Intrinsic::x86_avx512_permvar_hi_256;
2958 else if (VecWidth == 512 && EltWidth == 16)
2959 IID = Intrinsic::x86_avx512_permvar_hi_512;
2960 else if (VecWidth == 128 && EltWidth == 8)
2961 IID = Intrinsic::x86_avx512_permvar_qi_128;
2962 else if (VecWidth == 256 && EltWidth == 8)
2963 IID = Intrinsic::x86_avx512_permvar_qi_256;
2964 else if (VecWidth == 512 && EltWidth == 8)
2965 IID = Intrinsic::x86_avx512_permvar_qi_512;
2968 }
else if (Name.starts_with(
"dbpsadbw.")) {
2969 if (VecWidth == 128)
2970 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2971 else if (VecWidth == 256)
2972 IID = Intrinsic::x86_avx512_dbpsadbw_256;
2973 else if (VecWidth == 512)
2974 IID = Intrinsic::x86_avx512_dbpsadbw_512;
2977 }
else if (Name.starts_with(
"pmultishift.qb.")) {
2978 if (VecWidth == 128)
2979 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
2980 else if (VecWidth == 256)
2981 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
2982 else if (VecWidth == 512)
2983 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
2986 }
else if (Name.starts_with(
"conflict.")) {
2987 if (Name[9] ==
'd' && VecWidth == 128)
2988 IID = Intrinsic::x86_avx512_conflict_d_128;
2989 else if (Name[9] ==
'd' && VecWidth == 256)
2990 IID = Intrinsic::x86_avx512_conflict_d_256;
2991 else if (Name[9] ==
'd' && VecWidth == 512)
2992 IID = Intrinsic::x86_avx512_conflict_d_512;
2993 else if (Name[9] ==
'q' && VecWidth == 128)
2994 IID = Intrinsic::x86_avx512_conflict_q_128;
2995 else if (Name[9] ==
'q' && VecWidth == 256)
2996 IID = Intrinsic::x86_avx512_conflict_q_256;
2997 else if (Name[9] ==
'q' && VecWidth == 512)
2998 IID = Intrinsic::x86_avx512_conflict_q_512;
3001 }
else if (Name.starts_with(
"pavg.")) {
3002 if (Name[5] ==
'b' && VecWidth == 128)
3003 IID = Intrinsic::x86_sse2_pavg_b;
3004 else if (Name[5] ==
'b' && VecWidth == 256)
3005 IID = Intrinsic::x86_avx2_pavg_b;
3006 else if (Name[5] ==
'b' && VecWidth == 512)
3007 IID = Intrinsic::x86_avx512_pavg_b_512;
3008 else if (Name[5] ==
'w' && VecWidth == 128)
3009 IID = Intrinsic::x86_sse2_pavg_w;
3010 else if (Name[5] ==
'w' && VecWidth == 256)
3011 IID = Intrinsic::x86_avx2_pavg_w;
3012 else if (Name[5] ==
'w' && VecWidth == 512)
3013 IID = Intrinsic::x86_avx512_pavg_w_512;
3022 Rep = Builder.CreateIntrinsic(IID, Args);
3033 if (AsmStr->find(
"mov\tfp") == 0 &&
3034 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3035 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3036 AsmStr->replace(Pos, 1,
";");
3042 Value *Rep =
nullptr;
3044 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3046 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3047 {Arg, Builder.getTrue()},
3049 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3050 Type *Ty = (Name ==
"abs.bf16")
3054 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3055 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3056 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3057 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3058 : Intrinsic::nvvm_fabs;
3059 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3060 }
else if (Name.consume_front(
"ex2.approx.")) {
3062 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3063 : Intrinsic::nvvm_ex2_approx;
3064 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3065 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3066 Name.starts_with(
"atomic.load.add.f64.p")) {
3069 Rep = Builder.CreateAtomicRMW(
3075 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3076 Name.starts_with(
"atomic.load.dec.32.p")) {
3081 Rep = Builder.CreateAtomicRMW(
3085 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3091 Op.contains(
".cta.") ?
"block" :
"");
3092 if (
Op.starts_with(
"cas.")) {
3094 Value *Pair = Builder.CreateAtomicCmpXchg(
3097 Rep = Builder.CreateExtractValue(Pair, 0);
3115 "unexpected nvvm scoped atomic intrinsic");
3116 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3119 }
else if (Name ==
"clz.ll") {
3122 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3123 {Arg, Builder.getFalse()},
3125 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3126 }
else if (Name ==
"popc.ll") {
3130 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3131 Arg,
nullptr,
"ctpop");
3132 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3133 }
else if (Name ==
"h2f") {
3135 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3136 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3137 }
else if (Name.consume_front(
"bitcast.") &&
3138 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3141 }
else if (Name ==
"rotate.b32") {
3144 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3145 {Arg, Arg, ShiftAmt});
3146 }
else if (Name ==
"rotate.b64") {
3150 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3151 {Arg, Arg, ZExtShiftAmt});
3152 }
else if (Name ==
"rotate.right.b64") {
3156 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3157 {Arg, Arg, ZExtShiftAmt});
3158 }
else if (Name ==
"swap.lo.hi.b64") {
3161 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3162 {Arg, Arg, Builder.getInt64(32)});
3163 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3166 Name.starts_with(
".to.gen"))) {
3168 }
else if (Name.consume_front(
"ldg.global")) {
3172 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3175 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3177 }
else if (Name ==
"tanh.approx.f32") {
3181 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3183 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3185 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3186 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3188 }
else if (Name ==
"barrier") {
3189 Rep = Builder.CreateIntrinsic(
3190 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3192 }
else if (Name ==
"barrier.sync") {
3193 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3195 }
else if (Name ==
"barrier.sync.cnt") {
3196 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3198 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3199 Name ==
"barrier0.or") {
3201 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3205 .
Case(
"barrier0.popc",
3206 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3207 .
Case(
"barrier0.and",
3208 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3209 .
Case(
"barrier0.or",
3210 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3211 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3212 Rep = Builder.CreateZExt(Bar, CI->
getType());
3216 !
F->getReturnType()->getScalarType()->isBFloatTy()) {
3226 ? Builder.CreateBitCast(Arg, NewType)
3229 Rep = Builder.CreateCall(NewFn, Args);
3230 if (
F->getReturnType()->isIntegerTy())
3231 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3241 Value *Rep =
nullptr;
3243 if (Name.starts_with(
"sse4a.movnt.")) {
3255 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3258 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3259 }
else if (Name.starts_with(
"avx.movnt.") ||
3260 Name.starts_with(
"avx512.storent.")) {
3272 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3273 }
else if (Name ==
"sse2.storel.dq") {
3278 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3279 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3280 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3281 }
else if (Name.starts_with(
"sse.storeu.") ||
3282 Name.starts_with(
"sse2.storeu.") ||
3283 Name.starts_with(
"avx.storeu.")) {
3286 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3287 }
else if (Name ==
"avx512.mask.store.ss") {
3291 }
else if (Name.starts_with(
"avx512.mask.store")) {
3293 bool Aligned = Name[17] !=
'u';
3296 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3299 bool CmpEq = Name[9] ==
'e';
3302 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3303 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3310 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3311 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3313 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3314 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3315 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3316 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3317 Name.starts_with(
"sse2.sqrt.p") ||
3318 Name.starts_with(
"sse.sqrt.p")) {
3319 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3320 {CI->getArgOperand(0)});
3321 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3325 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3326 : Intrinsic::x86_avx512_sqrt_pd_512;
3329 Rep = Builder.CreateIntrinsic(IID, Args);
3331 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3332 {CI->getArgOperand(0)});
3336 }
else if (Name.starts_with(
"avx512.ptestm") ||
3337 Name.starts_with(
"avx512.ptestnm")) {
3341 Rep = Builder.CreateAnd(Op0, Op1);
3347 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3349 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3352 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3355 }
else if (Name.starts_with(
"avx512.kunpck")) {
3360 for (
unsigned i = 0; i != NumElts; ++i)
3369 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3370 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3371 }
else if (Name ==
"avx512.kand.w") {
3374 Rep = Builder.CreateAnd(
LHS,
RHS);
3375 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3376 }
else if (Name ==
"avx512.kandn.w") {
3379 LHS = Builder.CreateNot(
LHS);
3380 Rep = Builder.CreateAnd(
LHS,
RHS);
3381 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3382 }
else if (Name ==
"avx512.kor.w") {
3385 Rep = Builder.CreateOr(
LHS,
RHS);
3386 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3387 }
else if (Name ==
"avx512.kxor.w") {
3390 Rep = Builder.CreateXor(
LHS,
RHS);
3391 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3392 }
else if (Name ==
"avx512.kxnor.w") {
3395 LHS = Builder.CreateNot(
LHS);
3396 Rep = Builder.CreateXor(
LHS,
RHS);
3397 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3398 }
else if (Name ==
"avx512.knot.w") {
3400 Rep = Builder.CreateNot(Rep);
3401 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3402 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3405 Rep = Builder.CreateOr(
LHS,
RHS);
3406 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3408 if (Name[14] ==
'c')
3412 Rep = Builder.CreateICmpEQ(Rep,
C);
3413 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3414 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3415 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3416 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3417 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3420 ConstantInt::get(I32Ty, 0));
3422 ConstantInt::get(I32Ty, 0));
3424 if (Name.contains(
".add."))
3425 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3426 else if (Name.contains(
".sub."))
3427 EltOp = Builder.CreateFSub(Elt0, Elt1);
3428 else if (Name.contains(
".mul."))
3429 EltOp = Builder.CreateFMul(Elt0, Elt1);
3431 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3432 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3433 ConstantInt::get(I32Ty, 0));
3434 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3436 bool CmpEq = Name[16] ==
'e';
3438 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3440 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3447 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3450 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3453 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3460 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3462 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3463 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3465 if (VecWidth == 128 && EltWidth == 32)
3466 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3467 else if (VecWidth == 256 && EltWidth == 32)
3468 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3469 else if (VecWidth == 512 && EltWidth == 32)
3470 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3471 else if (VecWidth == 128 && EltWidth == 64)
3472 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3473 else if (VecWidth == 256 && EltWidth == 64)
3474 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3475 else if (VecWidth == 512 && EltWidth == 64)
3476 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3483 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3486 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3487 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3489 if (VecWidth == 128 && EltWidth == 32)
3490 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3491 else if (VecWidth == 256 && EltWidth == 32)
3492 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3493 else if (VecWidth == 512 && EltWidth == 32)
3494 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3495 else if (VecWidth == 128 && EltWidth == 64)
3496 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3497 else if (VecWidth == 256 && EltWidth == 64)
3498 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3499 else if (VecWidth == 512 && EltWidth == 64)
3500 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3505 if (VecWidth == 512)
3507 Args.push_back(Mask);
3509 Rep = Builder.CreateIntrinsic(IID, Args);
3510 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3514 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3517 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3518 Name.starts_with(
"avx512.cvtw2mask.") ||
3519 Name.starts_with(
"avx512.cvtd2mask.") ||
3520 Name.starts_with(
"avx512.cvtq2mask.")) {
3525 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3526 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3527 Name.starts_with(
"avx512.mask.pabs")) {
3529 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3530 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3531 Name.starts_with(
"avx512.mask.pmaxs")) {
3533 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3534 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3535 Name.starts_with(
"avx512.mask.pmaxu")) {
3537 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3538 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3539 Name.starts_with(
"avx512.mask.pmins")) {
3541 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3542 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3543 Name.starts_with(
"avx512.mask.pminu")) {
3545 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3546 Name ==
"avx512.pmulu.dq.512" ||
3547 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3549 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3550 Name ==
"avx512.pmul.dq.512" ||
3551 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3553 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3554 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3559 }
else if (Name ==
"avx512.cvtusi2sd") {
3564 }
else if (Name ==
"sse2.cvtss2sd") {
3566 Rep = Builder.CreateFPExt(
3569 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3570 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3571 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3572 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3573 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3574 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3575 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3576 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3577 Name ==
"avx512.mask.cvtqq2ps.256" ||
3578 Name ==
"avx512.mask.cvtqq2ps.512" ||
3579 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3580 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3581 Name ==
"avx.cvt.ps2.pd.256" ||
3582 Name ==
"avx512.mask.cvtps2pd.128" ||
3583 Name ==
"avx512.mask.cvtps2pd.256") {
3588 unsigned NumDstElts = DstTy->getNumElements();
3589 if (NumDstElts < SrcTy->getNumElements()) {
3590 assert(NumDstElts == 2 &&
"Unexpected vector size");
3591 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3594 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3595 bool IsUnsigned = Name.contains(
"cvtu");
3597 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3601 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3602 : Intrinsic::x86_avx512_sitofp_round;
3603 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3606 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3607 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3613 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3614 Name.starts_with(
"vcvtph2ps.")) {
3618 unsigned NumDstElts = DstTy->getNumElements();
3619 if (NumDstElts != SrcTy->getNumElements()) {
3620 assert(NumDstElts == 4 &&
"Unexpected vector size");
3621 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3623 Rep = Builder.CreateBitCast(
3625 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3629 }
else if (Name.starts_with(
"avx512.mask.load")) {
3631 bool Aligned = Name[16] !=
'u';
3634 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3638 ResultTy->getNumElements());
3639 Rep = Builder.CreateIntrinsic(
3640 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3642 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3648 Rep = Builder.CreateIntrinsic(
3649 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3651 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3652 Name.starts_with(
"avx512.mask.expand.")) {
3656 ResultTy->getNumElements());
3658 bool IsCompress = Name[12] ==
'c';
3659 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3660 : Intrinsic::x86_avx512_mask_expand;
3661 Rep = Builder.CreateIntrinsic(
3663 }
else if (Name.starts_with(
"xop.vpcom")) {
3665 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3666 Name.ends_with(
"uq"))
3668 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3669 Name.ends_with(
"d") || Name.ends_with(
"q"))
3678 Name = Name.substr(9);
3679 if (Name.starts_with(
"lt"))
3681 else if (Name.starts_with(
"le"))
3683 else if (Name.starts_with(
"gt"))
3685 else if (Name.starts_with(
"ge"))
3687 else if (Name.starts_with(
"eq"))
3689 else if (Name.starts_with(
"ne"))
3691 else if (Name.starts_with(
"false"))
3693 else if (Name.starts_with(
"true"))
3700 }
else if (Name.starts_with(
"xop.vpcmov")) {
3702 Value *NotSel = Builder.CreateNot(Sel);
3705 Rep = Builder.CreateOr(Sel0, Sel1);
3706 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3707 Name.starts_with(
"avx512.mask.prol")) {
3709 }
else if (Name.starts_with(
"avx512.pror") ||
3710 Name.starts_with(
"avx512.mask.pror")) {
3712 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3713 Name.starts_with(
"avx512.mask.vpshld") ||
3714 Name.starts_with(
"avx512.maskz.vpshld")) {
3715 bool ZeroMask = Name[11] ==
'z';
3717 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3718 Name.starts_with(
"avx512.mask.vpshrd") ||
3719 Name.starts_with(
"avx512.maskz.vpshrd")) {
3720 bool ZeroMask = Name[11] ==
'z';
3722 }
else if (Name ==
"sse42.crc32.64.8") {
3725 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3727 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3728 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3729 Name.starts_with(
"avx512.vbroadcast.s")) {
3732 Type *EltTy = VecTy->getElementType();
3733 unsigned EltNum = VecTy->getNumElements();
3737 for (
unsigned I = 0;
I < EltNum; ++
I)
3738 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3739 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3740 Name.starts_with(
"sse41.pmovzx") ||
3741 Name.starts_with(
"avx2.pmovsx") ||
3742 Name.starts_with(
"avx2.pmovzx") ||
3743 Name.starts_with(
"avx512.mask.pmovsx") ||
3744 Name.starts_with(
"avx512.mask.pmovzx")) {
3746 unsigned NumDstElts = DstTy->getNumElements();
3750 for (
unsigned i = 0; i != NumDstElts; ++i)
3755 bool DoSext = Name.contains(
"pmovsx");
3757 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3762 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3763 Name ==
"avx512.mask.pmov.qd.512" ||
3764 Name ==
"avx512.mask.pmov.wb.256" ||
3765 Name ==
"avx512.mask.pmov.wb.512") {
3770 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3771 Name ==
"avx2.vbroadcasti128") {
3777 if (NumSrcElts == 2)
3780 Rep = Builder.CreateShuffleVector(
Load,
3782 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3783 Name.starts_with(
"avx512.mask.shuf.f")) {
3788 unsigned ControlBitsMask = NumLanes - 1;
3789 unsigned NumControlBits = NumLanes / 2;
3792 for (
unsigned l = 0; l != NumLanes; ++l) {
3793 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3795 if (l >= NumLanes / 2)
3796 LaneMask += NumLanes;
3797 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3798 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3804 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3805 Name.starts_with(
"avx512.mask.broadcasti")) {
3808 unsigned NumDstElts =
3812 for (
unsigned i = 0; i != NumDstElts; ++i)
3813 ShuffleMask[i] = i % NumSrcElts;
3819 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3820 Name.starts_with(
"avx2.vbroadcast") ||
3821 Name.starts_with(
"avx512.pbroadcast") ||
3822 Name.starts_with(
"avx512.mask.broadcast.s")) {
3829 Rep = Builder.CreateShuffleVector(
Op, M);
3834 }
else if (Name.starts_with(
"sse2.padds.") ||
3835 Name.starts_with(
"avx2.padds.") ||
3836 Name.starts_with(
"avx512.padds.") ||
3837 Name.starts_with(
"avx512.mask.padds.")) {
3839 }
else if (Name.starts_with(
"sse2.psubs.") ||
3840 Name.starts_with(
"avx2.psubs.") ||
3841 Name.starts_with(
"avx512.psubs.") ||
3842 Name.starts_with(
"avx512.mask.psubs.")) {
3844 }
else if (Name.starts_with(
"sse2.paddus.") ||
3845 Name.starts_with(
"avx2.paddus.") ||
3846 Name.starts_with(
"avx512.mask.paddus.")) {
3848 }
else if (Name.starts_with(
"sse2.psubus.") ||
3849 Name.starts_with(
"avx2.psubus.") ||
3850 Name.starts_with(
"avx512.mask.psubus.")) {
3852 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3857 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3861 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3866 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3871 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3872 Name ==
"avx512.psll.dq.512") {
3876 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3877 Name ==
"avx512.psrl.dq.512") {
3881 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3882 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3883 Name.starts_with(
"avx2.pblendd.")) {
3888 unsigned NumElts = VecTy->getNumElements();
3891 for (
unsigned i = 0; i != NumElts; ++i)
3892 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
3894 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3895 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3896 Name ==
"avx2.vinserti128" ||
3897 Name.starts_with(
"avx512.mask.insert")) {
3901 unsigned DstNumElts =
3903 unsigned SrcNumElts =
3905 unsigned Scale = DstNumElts / SrcNumElts;
3912 for (
unsigned i = 0; i != SrcNumElts; ++i)
3914 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3915 Idxs[i] = SrcNumElts;
3916 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3930 for (
unsigned i = 0; i != DstNumElts; ++i)
3933 for (
unsigned i = 0; i != SrcNumElts; ++i)
3934 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
3935 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3941 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3942 Name ==
"avx2.vextracti128" ||
3943 Name.starts_with(
"avx512.mask.vextract")) {
3946 unsigned DstNumElts =
3948 unsigned SrcNumElts =
3950 unsigned Scale = SrcNumElts / DstNumElts;
3957 for (
unsigned i = 0; i != DstNumElts; ++i) {
3958 Idxs[i] = i + (
Imm * DstNumElts);
3960 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3966 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3967 Name.starts_with(
"avx512.mask.perm.di.")) {
3971 unsigned NumElts = VecTy->getNumElements();
3974 for (
unsigned i = 0; i != NumElts; ++i)
3975 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
3977 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3982 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
3994 unsigned HalfSize = NumElts / 2;
4006 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4007 for (
unsigned i = 0; i < HalfSize; ++i)
4008 ShuffleMask[i] = StartIndex + i;
4011 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4012 for (
unsigned i = 0; i < HalfSize; ++i)
4013 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4015 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4017 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4018 Name.starts_with(
"avx512.mask.vpermil.p") ||
4019 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4023 unsigned NumElts = VecTy->getNumElements();
4025 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4026 unsigned IdxMask = ((1 << IdxSize) - 1);
4032 for (
unsigned i = 0; i != NumElts; ++i)
4033 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4035 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4040 }
else if (Name ==
"sse2.pshufl.w" ||
4041 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4046 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4050 for (
unsigned l = 0; l != NumElts; l += 8) {
4051 for (
unsigned i = 0; i != 4; ++i)
4052 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4053 for (
unsigned i = 4; i != 8; ++i)
4054 Idxs[i + l] = i + l;
4057 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4062 }
else if (Name ==
"sse2.pshufh.w" ||
4063 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4068 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4072 for (
unsigned l = 0; l != NumElts; l += 8) {
4073 for (
unsigned i = 0; i != 4; ++i)
4074 Idxs[i + l] = i + l;
4075 for (
unsigned i = 0; i != 4; ++i)
4076 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4079 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4084 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4091 unsigned HalfLaneElts = NumLaneElts / 2;
4094 for (
unsigned i = 0; i != NumElts; ++i) {
4096 Idxs[i] = i - (i % NumLaneElts);
4098 if ((i % NumLaneElts) >= HalfLaneElts)
4102 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4105 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4109 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4110 Name.starts_with(
"avx512.mask.movshdup") ||
4111 Name.starts_with(
"avx512.mask.movsldup")) {
4117 if (Name.starts_with(
"avx512.mask.movshdup."))
4121 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4122 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4123 Idxs[i + l + 0] = i + l +
Offset;
4124 Idxs[i + l + 1] = i + l +
Offset;
4127 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4131 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4132 Name.starts_with(
"avx512.mask.unpckl.")) {
4139 for (
int l = 0; l != NumElts; l += NumLaneElts)
4140 for (
int i = 0; i != NumLaneElts; ++i)
4141 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4143 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4147 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4148 Name.starts_with(
"avx512.mask.unpckh.")) {
4155 for (
int l = 0; l != NumElts; l += NumLaneElts)
4156 for (
int i = 0; i != NumLaneElts; ++i)
4157 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4159 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4163 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4164 Name.starts_with(
"avx512.mask.pand.")) {
4167 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4169 Rep = Builder.CreateBitCast(Rep, FTy);
4172 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4173 Name.starts_with(
"avx512.mask.pandn.")) {
4176 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4177 Rep = Builder.CreateAnd(Rep,
4179 Rep = Builder.CreateBitCast(Rep, FTy);
4182 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4183 Name.starts_with(
"avx512.mask.por.")) {
4186 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4188 Rep = Builder.CreateBitCast(Rep, FTy);
4191 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4192 Name.starts_with(
"avx512.mask.pxor.")) {
4195 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4197 Rep = Builder.CreateBitCast(Rep, FTy);
4200 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4204 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4208 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4212 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4213 if (Name.ends_with(
".512")) {
4215 if (Name[17] ==
's')
4216 IID = Intrinsic::x86_avx512_add_ps_512;
4218 IID = Intrinsic::x86_avx512_add_pd_512;
4220 Rep = Builder.CreateIntrinsic(
4228 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4229 if (Name.ends_with(
".512")) {
4231 if (Name[17] ==
's')
4232 IID = Intrinsic::x86_avx512_div_ps_512;
4234 IID = Intrinsic::x86_avx512_div_pd_512;
4236 Rep = Builder.CreateIntrinsic(
4244 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4245 if (Name.ends_with(
".512")) {
4247 if (Name[17] ==
's')
4248 IID = Intrinsic::x86_avx512_mul_ps_512;
4250 IID = Intrinsic::x86_avx512_mul_pd_512;
4252 Rep = Builder.CreateIntrinsic(
4260 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4261 if (Name.ends_with(
".512")) {
4263 if (Name[17] ==
's')
4264 IID = Intrinsic::x86_avx512_sub_ps_512;
4266 IID = Intrinsic::x86_avx512_sub_pd_512;
4268 Rep = Builder.CreateIntrinsic(
4276 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4277 Name.starts_with(
"avx512.mask.min.p")) &&
4278 Name.drop_front(18) ==
".512") {
4279 bool IsDouble = Name[17] ==
'd';
4280 bool IsMin = Name[13] ==
'i';
4282 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4283 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4286 Rep = Builder.CreateIntrinsic(
4291 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4293 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4294 {CI->getArgOperand(0), Builder.getInt1(false)});
4297 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4298 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4299 bool IsVariable = Name[16] ==
'v';
4300 char Size = Name[16] ==
'.' ? Name[17]
4301 : Name[17] ==
'.' ? Name[18]
4302 : Name[18] ==
'.' ? Name[19]
4306 if (IsVariable && Name[17] !=
'.') {
4307 if (
Size ==
'd' && Name[17] ==
'2')
4308 IID = Intrinsic::x86_avx2_psllv_q;
4309 else if (
Size ==
'd' && Name[17] ==
'4')
4310 IID = Intrinsic::x86_avx2_psllv_q_256;
4311 else if (
Size ==
's' && Name[17] ==
'4')
4312 IID = Intrinsic::x86_avx2_psllv_d;
4313 else if (
Size ==
's' && Name[17] ==
'8')
4314 IID = Intrinsic::x86_avx2_psllv_d_256;
4315 else if (
Size ==
'h' && Name[17] ==
'8')
4316 IID = Intrinsic::x86_avx512_psllv_w_128;
4317 else if (
Size ==
'h' && Name[17] ==
'1')
4318 IID = Intrinsic::x86_avx512_psllv_w_256;
4319 else if (Name[17] ==
'3' && Name[18] ==
'2')
4320 IID = Intrinsic::x86_avx512_psllv_w_512;
4323 }
else if (Name.ends_with(
".128")) {
4325 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4326 : Intrinsic::x86_sse2_psll_d;
4327 else if (
Size ==
'q')
4328 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4329 : Intrinsic::x86_sse2_psll_q;
4330 else if (
Size ==
'w')
4331 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4332 : Intrinsic::x86_sse2_psll_w;
4335 }
else if (Name.ends_with(
".256")) {
4337 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4338 : Intrinsic::x86_avx2_psll_d;
4339 else if (
Size ==
'q')
4340 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4341 : Intrinsic::x86_avx2_psll_q;
4342 else if (
Size ==
'w')
4343 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4344 : Intrinsic::x86_avx2_psll_w;
4349 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4350 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4351 : Intrinsic::x86_avx512_psll_d_512;
4352 else if (
Size ==
'q')
4353 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4354 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4355 : Intrinsic::x86_avx512_psll_q_512;
4356 else if (
Size ==
'w')
4357 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4358 : Intrinsic::x86_avx512_psll_w_512;
4364 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4365 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4366 bool IsVariable = Name[16] ==
'v';
4367 char Size = Name[16] ==
'.' ? Name[17]
4368 : Name[17] ==
'.' ? Name[18]
4369 : Name[18] ==
'.' ? Name[19]
4373 if (IsVariable && Name[17] !=
'.') {
4374 if (
Size ==
'd' && Name[17] ==
'2')
4375 IID = Intrinsic::x86_avx2_psrlv_q;
4376 else if (
Size ==
'd' && Name[17] ==
'4')
4377 IID = Intrinsic::x86_avx2_psrlv_q_256;
4378 else if (
Size ==
's' && Name[17] ==
'4')
4379 IID = Intrinsic::x86_avx2_psrlv_d;
4380 else if (
Size ==
's' && Name[17] ==
'8')
4381 IID = Intrinsic::x86_avx2_psrlv_d_256;
4382 else if (
Size ==
'h' && Name[17] ==
'8')
4383 IID = Intrinsic::x86_avx512_psrlv_w_128;
4384 else if (
Size ==
'h' && Name[17] ==
'1')
4385 IID = Intrinsic::x86_avx512_psrlv_w_256;
4386 else if (Name[17] ==
'3' && Name[18] ==
'2')
4387 IID = Intrinsic::x86_avx512_psrlv_w_512;
4390 }
else if (Name.ends_with(
".128")) {
4392 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4393 : Intrinsic::x86_sse2_psrl_d;
4394 else if (
Size ==
'q')
4395 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4396 : Intrinsic::x86_sse2_psrl_q;
4397 else if (
Size ==
'w')
4398 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4399 : Intrinsic::x86_sse2_psrl_w;
4402 }
else if (Name.ends_with(
".256")) {
4404 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4405 : Intrinsic::x86_avx2_psrl_d;
4406 else if (
Size ==
'q')
4407 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4408 : Intrinsic::x86_avx2_psrl_q;
4409 else if (
Size ==
'w')
4410 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4411 : Intrinsic::x86_avx2_psrl_w;
4416 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4417 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4418 : Intrinsic::x86_avx512_psrl_d_512;
4419 else if (
Size ==
'q')
4420 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4421 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4422 : Intrinsic::x86_avx512_psrl_q_512;
4423 else if (
Size ==
'w')
4424 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4425 : Intrinsic::x86_avx512_psrl_w_512;
4431 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4432 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4433 bool IsVariable = Name[16] ==
'v';
4434 char Size = Name[16] ==
'.' ? Name[17]
4435 : Name[17] ==
'.' ? Name[18]
4436 : Name[18] ==
'.' ? Name[19]
4440 if (IsVariable && Name[17] !=
'.') {
4441 if (
Size ==
's' && Name[17] ==
'4')
4442 IID = Intrinsic::x86_avx2_psrav_d;
4443 else if (
Size ==
's' && Name[17] ==
'8')
4444 IID = Intrinsic::x86_avx2_psrav_d_256;
4445 else if (
Size ==
'h' && Name[17] ==
'8')
4446 IID = Intrinsic::x86_avx512_psrav_w_128;
4447 else if (
Size ==
'h' && Name[17] ==
'1')
4448 IID = Intrinsic::x86_avx512_psrav_w_256;
4449 else if (Name[17] ==
'3' && Name[18] ==
'2')
4450 IID = Intrinsic::x86_avx512_psrav_w_512;
4453 }
else if (Name.ends_with(
".128")) {
4455 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4456 : Intrinsic::x86_sse2_psra_d;
4457 else if (
Size ==
'q')
4458 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4459 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4460 : Intrinsic::x86_avx512_psra_q_128;
4461 else if (
Size ==
'w')
4462 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4463 : Intrinsic::x86_sse2_psra_w;
4466 }
else if (Name.ends_with(
".256")) {
4468 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4469 : Intrinsic::x86_avx2_psra_d;
4470 else if (
Size ==
'q')
4471 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4472 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4473 : Intrinsic::x86_avx512_psra_q_256;
4474 else if (
Size ==
'w')
4475 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4476 : Intrinsic::x86_avx2_psra_w;
4481 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4482 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4483 : Intrinsic::x86_avx512_psra_d_512;
4484 else if (
Size ==
'q')
4485 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4486 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4487 : Intrinsic::x86_avx512_psra_q_512;
4488 else if (
Size ==
'w')
4489 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4490 : Intrinsic::x86_avx512_psra_w_512;
4496 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4498 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4500 }
else if (Name.ends_with(
".movntdqa")) {
4504 LoadInst *LI = Builder.CreateAlignedLoad(
4509 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4510 Name.starts_with(
"fma.vfmsub.") ||
4511 Name.starts_with(
"fma.vfnmadd.") ||
4512 Name.starts_with(
"fma.vfnmsub.")) {
4513 bool NegMul = Name[6] ==
'n';
4514 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4515 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4526 if (NegMul && !IsScalar)
4527 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4528 if (NegMul && IsScalar)
4529 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4531 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4533 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4537 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4545 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4549 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4550 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4551 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4552 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4553 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4554 bool IsMask3 = Name[11] ==
'3';
4555 bool IsMaskZ = Name[11] ==
'z';
4557 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4558 bool NegMul = Name[2] ==
'n';
4559 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4565 if (NegMul && (IsMask3 || IsMaskZ))
4566 A = Builder.CreateFNeg(
A);
4567 if (NegMul && !(IsMask3 || IsMaskZ))
4568 B = Builder.CreateFNeg(
B);
4570 C = Builder.CreateFNeg(
C);
4572 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4573 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4574 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4581 if (Name.back() ==
'd')
4582 IID = Intrinsic::x86_avx512_vfmadd_f64;
4584 IID = Intrinsic::x86_avx512_vfmadd_f32;
4585 Rep = Builder.CreateIntrinsic(IID,
Ops);
4587 Rep = Builder.CreateFMA(
A,
B,
C);
4596 if (NegAcc && IsMask3)
4601 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4603 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4604 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4605 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4606 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4607 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4608 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4609 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4610 bool IsMask3 = Name[11] ==
'3';
4611 bool IsMaskZ = Name[11] ==
'z';
4613 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4614 bool NegMul = Name[2] ==
'n';
4615 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4621 if (NegMul && (IsMask3 || IsMaskZ))
4622 A = Builder.CreateFNeg(
A);
4623 if (NegMul && !(IsMask3 || IsMaskZ))
4624 B = Builder.CreateFNeg(
B);
4626 C = Builder.CreateFNeg(
C);
4633 if (Name[Name.size() - 5] ==
's')
4634 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4636 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4640 Rep = Builder.CreateFMA(
A,
B,
C);
4648 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4652 if (VecWidth == 128 && EltWidth == 32)
4653 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4654 else if (VecWidth == 256 && EltWidth == 32)
4655 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4656 else if (VecWidth == 128 && EltWidth == 64)
4657 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4658 else if (VecWidth == 256 && EltWidth == 64)
4659 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4665 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4666 Rep = Builder.CreateIntrinsic(IID,
Ops);
4667 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4668 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4669 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4670 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4671 bool IsMask3 = Name[11] ==
'3';
4672 bool IsMaskZ = Name[11] ==
'z';
4674 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4675 bool IsSubAdd = Name[3] ==
's';
4679 if (Name[Name.size() - 5] ==
's')
4680 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4682 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4687 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4689 Rep = Builder.CreateIntrinsic(IID,
Ops);
4698 Value *Odd = Builder.CreateCall(FMA,
Ops);
4699 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4700 Value *Even = Builder.CreateCall(FMA,
Ops);
4706 for (
int i = 0; i != NumElts; ++i)
4707 Idxs[i] = i + (i % 2) * NumElts;
4709 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4717 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4718 Name.starts_with(
"avx512.maskz.pternlog.")) {
4719 bool ZeroMask = Name[11] ==
'z';
4723 if (VecWidth == 128 && EltWidth == 32)
4724 IID = Intrinsic::x86_avx512_pternlog_d_128;
4725 else if (VecWidth == 256 && EltWidth == 32)
4726 IID = Intrinsic::x86_avx512_pternlog_d_256;
4727 else if (VecWidth == 512 && EltWidth == 32)
4728 IID = Intrinsic::x86_avx512_pternlog_d_512;
4729 else if (VecWidth == 128 && EltWidth == 64)
4730 IID = Intrinsic::x86_avx512_pternlog_q_128;
4731 else if (VecWidth == 256 && EltWidth == 64)
4732 IID = Intrinsic::x86_avx512_pternlog_q_256;
4733 else if (VecWidth == 512 && EltWidth == 64)
4734 IID = Intrinsic::x86_avx512_pternlog_q_512;
4740 Rep = Builder.CreateIntrinsic(IID, Args);
4744 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4745 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4746 bool ZeroMask = Name[11] ==
'z';
4747 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4750 if (VecWidth == 128 && !
High)
4751 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4752 else if (VecWidth == 256 && !
High)
4753 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4754 else if (VecWidth == 512 && !
High)
4755 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4756 else if (VecWidth == 128 &&
High)
4757 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4758 else if (VecWidth == 256 &&
High)
4759 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4760 else if (VecWidth == 512 &&
High)
4761 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4767 Rep = Builder.CreateIntrinsic(IID, Args);
4771 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4772 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4773 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4774 bool ZeroMask = Name[11] ==
'z';
4775 bool IndexForm = Name[17] ==
'i';
4777 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4778 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4779 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4780 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4781 bool ZeroMask = Name[11] ==
'z';
4782 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4785 if (VecWidth == 128 && !IsSaturating)
4786 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4787 else if (VecWidth == 256 && !IsSaturating)
4788 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4789 else if (VecWidth == 512 && !IsSaturating)
4790 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4791 else if (VecWidth == 128 && IsSaturating)
4792 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4793 else if (VecWidth == 256 && IsSaturating)
4794 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4795 else if (VecWidth == 512 && IsSaturating)
4796 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4806 if (Args[1]->
getType()->isVectorTy() &&
4809 ->isIntegerTy(32) &&
4810 Args[2]->
getType()->isVectorTy() &&
4813 ->isIntegerTy(32)) {
4814 Type *NewArgType =
nullptr;
4815 if (VecWidth == 128)
4817 else if (VecWidth == 256)
4819 else if (VecWidth == 512)
4825 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4826 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4829 Rep = Builder.CreateIntrinsic(IID, Args);
4833 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4834 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4835 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4836 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4837 bool ZeroMask = Name[11] ==
'z';
4838 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4841 if (VecWidth == 128 && !IsSaturating)
4842 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4843 else if (VecWidth == 256 && !IsSaturating)
4844 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4845 else if (VecWidth == 512 && !IsSaturating)
4846 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4847 else if (VecWidth == 128 && IsSaturating)
4848 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4849 else if (VecWidth == 256 && IsSaturating)
4850 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4851 else if (VecWidth == 512 && IsSaturating)
4852 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4862 if (Args[1]->
getType()->isVectorTy() &&
4865 ->isIntegerTy(32) &&
4866 Args[2]->
getType()->isVectorTy() &&
4869 ->isIntegerTy(32)) {
4870 Type *NewArgType =
nullptr;
4871 if (VecWidth == 128)
4873 else if (VecWidth == 256)
4875 else if (VecWidth == 512)
4881 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4882 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4885 Rep = Builder.CreateIntrinsic(IID, Args);
4889 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4890 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4891 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4893 if (Name[0] ==
'a' && Name.back() ==
'2')
4894 IID = Intrinsic::x86_addcarry_32;
4895 else if (Name[0] ==
'a' && Name.back() ==
'4')
4896 IID = Intrinsic::x86_addcarry_64;
4897 else if (Name[0] ==
's' && Name.back() ==
'2')
4898 IID = Intrinsic::x86_subborrow_32;
4899 else if (Name[0] ==
's' && Name.back() ==
'4')
4900 IID = Intrinsic::x86_subborrow_64;
4907 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4910 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4913 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4917 }
else if (Name.starts_with(
"avx512.mask.") &&
4920 }
else if (Name.starts_with(
"bmi.pdep.")) {
4922 }
else if (Name.starts_with(
"bmi.pext.")) {
4932 if (Name.starts_with(
"neon.bfcvt")) {
4933 if (Name.starts_with(
"neon.bfcvtn2")) {
4935 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4937 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4938 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4941 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4942 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4944 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4948 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4949 return Builder.CreateShuffleVector(
4952 return Builder.CreateFPTrunc(CI->
getOperand(0),
4955 }
else if (Name.starts_with(
"sve.fcvt")) {
4958 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4959 .
Case(
"sve.fcvtnt.bf16f32",
4960 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4972 if (Args[1]->
getType() != BadPredTy)
4975 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
4976 BadPredTy, Args[1]);
4977 Args[1] = Builder.CreateIntrinsic(
4978 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
4980 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
4984 if (Name ==
"neon.vcvtfp2hf")
4985 return Builder.CreateBitCast(
4986 Builder.CreateFPTrunc(
4990 if (Name ==
"neon.vcvthf2fp")
4991 return Builder.CreateFPExt(
4992 Builder.CreateBitCast(
5002 if (Name ==
"mve.vctp64.old") {
5005 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5008 Value *C1 = Builder.CreateIntrinsic(
5009 Intrinsic::arm_mve_pred_v2i,
5011 return Builder.CreateIntrinsic(
5012 Intrinsic::arm_mve_pred_i2v,
5014 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5015 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5016 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5017 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5019 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5020 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5021 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5022 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5024 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5025 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5026 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5027 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5028 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5029 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5030 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5031 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5032 std::vector<Type *> Tys;
5036 case Intrinsic::arm_mve_mull_int_predicated:
5037 case Intrinsic::arm_mve_vqdmull_predicated:
5038 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5041 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5042 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5043 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5047 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5051 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5055 case Intrinsic::arm_cde_vcx1q_predicated:
5056 case Intrinsic::arm_cde_vcx1qa_predicated:
5057 case Intrinsic::arm_cde_vcx2q_predicated:
5058 case Intrinsic::arm_cde_vcx2qa_predicated:
5059 case Intrinsic::arm_cde_vcx3q_predicated:
5060 case Intrinsic::arm_cde_vcx3qa_predicated:
5067 std::vector<Value *>
Ops;
5069 Type *Ty =
Op->getType();
5070 if (Ty->getScalarSizeInBits() == 1) {
5071 Value *C1 = Builder.CreateIntrinsic(
5072 Intrinsic::arm_mve_pred_v2i,
5074 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5079 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5094 auto UpgradeLegacyWMMAIUIntrinsicCall =
5099 Args.push_back(Builder.getFalse());
5103 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5110 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5114 NewCall->copyMetadata(*CI);
5118 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5119 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5120 "intrinsic should have 7 arguments");
5123 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5125 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5126 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5127 "intrinsic should have 8 arguments");
5132 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5135 switch (
F->getIntrinsicID()) {
5138 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5139 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5140 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5141 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5142 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5143 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5158 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5161 F->getParent(),
F->getIntrinsicID(), Overloads);
5166 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5170 NewCall->copyMetadata(*CI);
5171 NewCall->takeName(CI);
5193 if (NumOperands < 3)
5206 bool IsVolatile =
false;
5210 if (NumOperands > 3)
5215 if (NumOperands > 5) {
5217 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5231 if (VT->getElementType()->isIntegerTy(16)) {
5234 Val = Builder.CreateBitCast(Val, AsBF16);
5242 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5244 unsigned AddrSpace = PtrTy->getAddressSpace();
5247 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5249 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5254 MDNode *RangeNotPrivate =
5257 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5263 return Builder.CreateBitCast(RMW, RetTy);
5284 return MAV->getMetadata();
5293 if (Name ==
"label") {
5295 }
else if (Name ==
"assign") {
5302 }
else if (Name ==
"declare") {
5306 }
else if (Name ==
"addr") {
5316 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5317 }
else if (Name ==
"value") {
5320 unsigned ExprOp = 2;
5335 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5343 int64_t OffsetVal =
Offset->getSExtValue();
5344 return Builder.CreateIntrinsic(OffsetVal >= 0
5345 ? Intrinsic::vector_splice_left
5346 : Intrinsic::vector_splice_right,
5348 {CI->getArgOperand(0), CI->getArgOperand(1),
5349 Builder.getInt32(std::abs(OffsetVal))});
5354 if (Name.starts_with(
"to.fp16")) {
5356 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5357 return Builder.CreateBitCast(Cast, CI->
getType());
5360 if (Name.starts_with(
"from.fp16")) {
5362 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5363 return Builder.CreateFPExt(Cast, CI->
getType());
5422 else if (Opcode == Instruction::ICmp)
5425 else if (Opcode == Instruction::FCmp)
5428 else if (Opcode == Instruction::Select)
5433 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5445 if (Defaults.empty())
5448 unsigned OldArgCount = CI->
arg_size();
5449 unsigned NewArgCount = NewFn->
arg_size();
5453 if (OldArgCount >= NewArgCount)
5461 if (OldArgCount < FirstDefault)
5466 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5467 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5468 "missing argument outside the default range");
5469 Type *ParamTy = NewFT->getParamType(Idx);
5474 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5480 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5512 if (!Name.consume_front(
"llvm."))
5515 bool IsX86 = Name.consume_front(
"x86.");
5516 bool IsNVVM = Name.consume_front(
"nvvm.");
5517 bool IsAArch64 = Name.consume_front(
"aarch64.");
5518 bool IsARM = Name.consume_front(
"arm.");
5519 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5520 bool IsDbg = Name.consume_front(
"dbg.");
5522 (Name.consume_front(
"experimental.vector.splice") ||
5523 Name.consume_front(
"vector.splice")) &&
5524 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5525 Value *Rep =
nullptr;
5527 if (!IsX86 && Name ==
"stackprotectorcheck") {
5529 }
else if (IsNVVM) {
5533 }
else if (IsAArch64) {
5537 }
else if (IsAMDGCN) {
5541 }
else if (IsOldSplice) {
5543 }
else if (Name.consume_front(
"convert.")) {
5545 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5560 const auto &DefaultCase = [&]() ->
void {
5568 "Unknown function for CallBase upgrade and isn't just a name change");
5576 "Return type must have changed");
5577 assert(OldST->getNumElements() ==
5579 "Must have same number of elements");
5582 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5585 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5586 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5587 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5611 case Intrinsic::arm_neon_vst1:
5612 case Intrinsic::arm_neon_vst2:
5613 case Intrinsic::arm_neon_vst3:
5614 case Intrinsic::arm_neon_vst4:
5615 case Intrinsic::arm_neon_vst2lane:
5616 case Intrinsic::arm_neon_vst3lane:
5617 case Intrinsic::arm_neon_vst4lane: {
5619 NewCall = Builder.CreateCall(NewFn, Args);
5622 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5623 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5624 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5629 NewCall = Builder.CreateCall(NewFn, Args);
5632 case Intrinsic::aarch64_sve_ld3_sret:
5633 case Intrinsic::aarch64_sve_ld4_sret:
5634 case Intrinsic::aarch64_sve_ld2_sret: {
5642 Name = Name.substr(5);
5649 unsigned MinElts = RetTy->getMinNumElements() /
N;
5651 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5653 for (
unsigned I = 0;
I <
N;
I++) {
5654 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5655 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5661 case Intrinsic::coro_end_async:
5662 case Intrinsic::coro_end: {
5664 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5666 NewCall = Builder.CreateCall(NewFn, Args);
5671 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5672 Value *InRamp = Builder.CreateCall(IsInRamp);
5682 case Intrinsic::vector_extract: {
5684 Name = Name.substr(5);
5685 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5690 unsigned MinElts = RetTy->getMinNumElements();
5693 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5697 case Intrinsic::vector_insert: {
5699 Name = Name.substr(5);
5700 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5704 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5709 NewCall = Builder.CreateCall(
5713 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5719 assert(
N > 1 &&
"Create is expected to be between 2-4");
5722 unsigned MinElts = RetTy->getMinNumElements() /
N;
5723 for (
unsigned I = 0;
I <
N;
I++) {
5725 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5732 case Intrinsic::arm_neon_bfdot:
5733 case Intrinsic::arm_neon_bfmmla:
5734 case Intrinsic::arm_neon_bfmlalb:
5735 case Intrinsic::arm_neon_bfmlalt:
5736 case Intrinsic::aarch64_neon_bfdot:
5737 case Intrinsic::aarch64_neon_bfmmla:
5738 case Intrinsic::aarch64_neon_bfmlalb:
5739 case Intrinsic::aarch64_neon_bfmlalt: {
5742 "Mismatch between function args and call args");
5743 size_t OperandWidth =
5745 assert((OperandWidth == 64 || OperandWidth == 128) &&
5746 "Unexpected operand width");
5748 auto Iter = CI->
args().begin();
5749 Args.push_back(*Iter++);
5750 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5751 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5752 NewCall = Builder.CreateCall(NewFn, Args);
5756 case Intrinsic::bitreverse:
5757 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5760 case Intrinsic::ctlz:
5761 case Intrinsic::cttz: {
5768 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5772 case Intrinsic::objectsize: {
5773 Value *NullIsUnknownSize =
5777 NewCall = Builder.CreateCall(
5782 case Intrinsic::ctpop:
5783 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5785 case Intrinsic::dbg_value: {
5787 Name = Name.substr(5);
5789 if (Name.starts_with(
"dbg.addr")) {
5803 if (
Offset->isNullValue()) {
5804 NewCall = Builder.CreateCall(
5813 case Intrinsic::ptr_annotation:
5821 NewCall = Builder.CreateCall(
5830 case Intrinsic::var_annotation:
5837 NewCall = Builder.CreateCall(
5846 case Intrinsic::riscv_aes32dsi:
5847 case Intrinsic::riscv_aes32dsmi:
5848 case Intrinsic::riscv_aes32esi:
5849 case Intrinsic::riscv_aes32esmi:
5850 case Intrinsic::riscv_sm4ks:
5851 case Intrinsic::riscv_sm4ed: {
5861 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5862 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5868 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5869 Value *Res = NewCall;
5871 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5877 case Intrinsic::nvvm_mapa_shared_cluster: {
5881 Value *Res = NewCall;
5882 Res = Builder.CreateAddrSpaceCast(
5889 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5890 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5893 Args[0] = Builder.CreateAddrSpaceCast(
5896 NewCall = Builder.CreateCall(NewFn, Args);
5902 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5903 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5904 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5905 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5906 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5907 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5908 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5909 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5916 Args[0] = Builder.CreateAddrSpaceCast(
5925 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5927 NewCall = Builder.CreateCall(NewFn, Args);
5933 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
5934 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
5935 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
5936 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
5937 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
5938 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
5939 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
5940 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
5942 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
5946 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
5947 NewCall = Builder.CreateCall(NewFn, Args);
5950 case Intrinsic::nvvm_tcgen05_mma_shared:
5951 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5952 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5953 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
5954 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
5955 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
5956 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
5957 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
5958 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
5959 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
5960 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
5961 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
5962 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
5963 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
5964 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
5965 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
5966 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
5967 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
5968 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
5969 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
5970 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
5971 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
5972 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
5973 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
5974 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
5975 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
5976 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
5977 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
5978 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
5979 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
5980 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
5981 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
5982 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
5983 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
5984 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
5985 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
5986 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
5987 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
5988 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
5990 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
5991 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
5993 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
5994 case Intrinsic::nvvm_tcgen05_mma_tensor:
5995 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
5996 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
5997 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
5998 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
5999 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
6000 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
6001 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
6002 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
6003 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
6004 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
6005 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
6006 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
6007 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
6008 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
6010 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
6011 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
6013 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
6015 Args.push_back(Builder.getInt32(0));
6016 NewCall = Builder.CreateCall(NewFn, Args);
6019 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6020 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6021 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6022 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6025 Builder.getFalse()});
6027 case Intrinsic::riscv_sha256sig0:
6028 case Intrinsic::riscv_sha256sig1:
6029 case Intrinsic::riscv_sha256sum0:
6030 case Intrinsic::riscv_sha256sum1:
6031 case Intrinsic::riscv_sm3p0:
6032 case Intrinsic::riscv_sm3p1: {
6039 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6041 NewCall = Builder.CreateCall(NewFn, Arg);
6043 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6050 case Intrinsic::x86_xop_vfrcz_ss:
6051 case Intrinsic::x86_xop_vfrcz_sd:
6052 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6055 case Intrinsic::x86_xop_vpermil2pd:
6056 case Intrinsic::x86_xop_vpermil2ps:
6057 case Intrinsic::x86_xop_vpermil2pd_256:
6058 case Intrinsic::x86_xop_vpermil2ps_256: {
6062 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6063 NewCall = Builder.CreateCall(NewFn, Args);
6067 case Intrinsic::x86_sse41_ptestc:
6068 case Intrinsic::x86_sse41_ptestz:
6069 case Intrinsic::x86_sse41_ptestnzc: {
6083 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6084 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6086 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6090 case Intrinsic::x86_rdtscp: {
6096 NewCall = Builder.CreateCall(NewFn);
6098 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6101 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6109 case Intrinsic::x86_sse41_insertps:
6110 case Intrinsic::x86_sse41_dppd:
6111 case Intrinsic::x86_sse41_dpps:
6112 case Intrinsic::x86_sse41_mpsadbw:
6113 case Intrinsic::x86_avx_dp_ps_256:
6114 case Intrinsic::x86_avx2_mpsadbw: {
6120 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6121 NewCall = Builder.CreateCall(NewFn, Args);
6125 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6126 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6127 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6128 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6129 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6130 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6136 NewCall = Builder.CreateCall(NewFn, Args);
6145 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6146 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6147 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6148 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6149 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6150 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6154 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6155 Args[1] = Builder.CreateBitCast(
6158 NewCall = Builder.CreateCall(NewFn, Args);
6159 Value *Res = Builder.CreateBitCast(
6167 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6168 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6169 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6173 Args[1] = Builder.CreateBitCast(
6175 Args[2] = Builder.CreateBitCast(
6178 NewCall = Builder.CreateCall(NewFn, Args);
6182 case Intrinsic::thread_pointer: {
6183 NewCall = Builder.CreateCall(NewFn, {});
6187 case Intrinsic::memcpy:
6188 case Intrinsic::memmove:
6189 case Intrinsic::memset: {
6205 NewCall = Builder.CreateCall(NewFn, Args);
6207 AttributeList NewAttrs = AttributeList::get(
6208 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6209 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6210 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6215 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6218 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6222 case Intrinsic::masked_load:
6223 case Intrinsic::masked_gather:
6224 case Intrinsic::masked_store:
6225 case Intrinsic::masked_scatter: {
6231 auto GetMaybeAlign = [](
Value *
Op) {
6233 uint64_t Val = CI->getZExtValue();
6241 auto GetAlign = [&](
Value *
Op) {
6250 case Intrinsic::masked_load:
6251 NewCall = Builder.CreateMaskedLoad(
6255 case Intrinsic::masked_gather:
6256 NewCall = Builder.CreateMaskedGather(
6262 case Intrinsic::masked_store:
6263 NewCall = Builder.CreateMaskedStore(
6267 case Intrinsic::masked_scatter:
6268 NewCall = Builder.CreateMaskedScatter(
6270 DL.getValueOrABITypeAlignment(
6284 case Intrinsic::lifetime_start:
6285 case Intrinsic::lifetime_end: {
6297 NewCall = Builder.CreateLifetimeStart(Ptr);
6299 NewCall = Builder.CreateLifetimeEnd(Ptr);
6308 case Intrinsic::x86_avx512_vpdpbusd_128:
6309 case Intrinsic::x86_avx512_vpdpbusd_256:
6310 case Intrinsic::x86_avx512_vpdpbusd_512:
6311 case Intrinsic::x86_avx512_vpdpbusds_128:
6312 case Intrinsic::x86_avx512_vpdpbusds_256:
6313 case Intrinsic::x86_avx512_vpdpbusds_512:
6314 case Intrinsic::x86_avx2_vpdpbssd_128:
6315 case Intrinsic::x86_avx2_vpdpbssd_256:
6316 case Intrinsic::x86_avx10_vpdpbssd_512:
6317 case Intrinsic::x86_avx2_vpdpbssds_128:
6318 case Intrinsic::x86_avx2_vpdpbssds_256:
6319 case Intrinsic::x86_avx10_vpdpbssds_512:
6320 case Intrinsic::x86_avx2_vpdpbsud_128:
6321 case Intrinsic::x86_avx2_vpdpbsud_256:
6322 case Intrinsic::x86_avx10_vpdpbsud_512:
6323 case Intrinsic::x86_avx2_vpdpbsuds_128:
6324 case Intrinsic::x86_avx2_vpdpbsuds_256:
6325 case Intrinsic::x86_avx10_vpdpbsuds_512:
6326 case Intrinsic::x86_avx2_vpdpbuud_128:
6327 case Intrinsic::x86_avx2_vpdpbuud_256:
6328 case Intrinsic::x86_avx10_vpdpbuud_512:
6329 case Intrinsic::x86_avx2_vpdpbuuds_128:
6330 case Intrinsic::x86_avx2_vpdpbuuds_256:
6331 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6336 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6337 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6339 NewCall = Builder.CreateCall(NewFn, Args);
6342 case Intrinsic::x86_avx512_vpdpwssd_128:
6343 case Intrinsic::x86_avx512_vpdpwssd_256:
6344 case Intrinsic::x86_avx512_vpdpwssd_512:
6345 case Intrinsic::x86_avx512_vpdpwssds_128:
6346 case Intrinsic::x86_avx512_vpdpwssds_256:
6347 case Intrinsic::x86_avx512_vpdpwssds_512:
6348 case Intrinsic::x86_avx2_vpdpwsud_128:
6349 case Intrinsic::x86_avx2_vpdpwsud_256:
6350 case Intrinsic::x86_avx10_vpdpwsud_512:
6351 case Intrinsic::x86_avx2_vpdpwsuds_128:
6352 case Intrinsic::x86_avx2_vpdpwsuds_256:
6353 case Intrinsic::x86_avx10_vpdpwsuds_512:
6354 case Intrinsic::x86_avx2_vpdpwusd_128:
6355 case Intrinsic::x86_avx2_vpdpwusd_256:
6356 case Intrinsic::x86_avx10_vpdpwusd_512:
6357 case Intrinsic::x86_avx2_vpdpwusds_128:
6358 case Intrinsic::x86_avx2_vpdpwusds_256:
6359 case Intrinsic::x86_avx10_vpdpwusds_512:
6360 case Intrinsic::x86_avx2_vpdpwuud_128:
6361 case Intrinsic::x86_avx2_vpdpwuud_256:
6362 case Intrinsic::x86_avx10_vpdpwuud_512:
6363 case Intrinsic::x86_avx2_vpdpwuuds_128:
6364 case Intrinsic::x86_avx2_vpdpwuuds_256:
6365 case Intrinsic::x86_avx10_vpdpwuuds_512:
6370 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6371 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6373 NewCall = Builder.CreateCall(NewFn, Args);
6376 assert(NewCall &&
"Should have either set this variable or returned through "
6377 "the default case");
6384 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6398 F->eraseFromParent();
6404 if (NumOperands == 0)
6412 if (NumOperands == 3) {
6416 Metadata *Elts2[] = {ScalarType, ScalarType,
6430 if (
Opc != Instruction::BitCast)
6434 Type *SrcTy = V->getType();
6451 if (
Opc != Instruction::BitCast)
6454 Type *SrcTy =
C->getType();
6471 if (Flag.getNumOperands() < 3)
6472 return std::nullopt;
6474 return Name->getString();
6475 return std::nullopt;
6489 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6490 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6491 if (auto Name = getModuleFlagNameSafely(*Flag))
6492 return *Name ==
"Debug Info Version";
6495 if (OpIt != ModFlags->op_end()) {
6496 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6503 bool BrokenDebugInfo =
false;
6506 if (!BrokenDebugInfo)
6512 M.getContext().diagnose(Diag);
6519 M.getContext().diagnose(DiagVersion);
6529 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6532 if (
F->hasFnAttribute(Attr)) {
6535 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6537 auto [Part, Rest] = S.
split(
',');
6543 const unsigned Dim = DimC -
'x';
6544 assert(Dim < 3 &&
"Unexpected dim char");
6554 F->addFnAttr(Attr, NewAttr);
6558 return S ==
"x" || S ==
"y" || S ==
"z";
6563 if (K ==
"kernel") {
6575 const unsigned Idx = (AlignIdxValuePair >> 16);
6576 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6581 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6586 if (K ==
"minctasm") {
6591 if (K ==
"maxnreg") {
6596 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6600 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6604 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6608 if (K ==
"grid_constant") {
6623 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6630 if (!SeenNodes.
insert(MD).second)
6637 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6644 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6646 const MDOperand &V = MD->getOperand(j + 1);
6649 NewOperands.
append({K, V});
6652 if (NewOperands.
size() > 1)
6665 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6666 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6667 if (ModRetainReleaseMarker) {
6673 ID->getString().split(ValueComp,
"#");
6674 if (ValueComp.
size() == 2) {
6675 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6679 M.eraseNamedMetadata(ModRetainReleaseMarker);
6690 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6716 bool InvalidCast =
false;
6718 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6731 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6733 Args.push_back(Arg);
6740 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6745 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6758 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6766 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6767 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6768 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6769 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6770 {
"objc_autoreleaseReturnValue",
6771 llvm::Intrinsic::objc_autoreleaseReturnValue},
6772 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6773 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6774 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6775 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6776 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6777 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6778 {
"objc_release", llvm::Intrinsic::objc_release},
6779 {
"objc_retain", llvm::Intrinsic::objc_retain},
6780 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6781 {
"objc_retainAutoreleaseReturnValue",
6782 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6783 {
"objc_retainAutoreleasedReturnValue",
6784 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6785 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6786 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6787 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6788 {
"objc_unsafeClaimAutoreleasedReturnValue",
6789 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6790 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6791 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6792 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6793 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6794 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6795 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6796 {
"objc_arc_annotation_topdown_bbstart",
6797 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6798 {
"objc_arc_annotation_topdown_bbend",
6799 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6800 {
"objc_arc_annotation_bottomup_bbstart",
6801 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6802 {
"objc_arc_annotation_bottomup_bbend",
6803 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6805 for (
auto &
I : RuntimeFuncs)
6806 UpgradeToIntrinsic(
I.first,
I.second);
6830 std::optional<bool> UseAddressDisc;
6833 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6834 for (
const MDNode *Flag : ModFlags->operands()) {
6836 if (Name && (*Name ==
"ptrauth-init-fini" ||
6837 *Name ==
"ptrauth-init-fini-address-discrimination"))
6842 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6843 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6844 constexpr unsigned ExpectedAddressMarker = 1;
6847 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6850 bool HasAddressDisc;
6851 if (!CPA->hasAddressDiscriminator())
6852 HasAddressDisc =
false;
6853 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6854 HasAddressDisc =
true;
6858 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6861 UseAddressDisc = HasAddressDisc;
6862 return CPA->getPointer();
6866 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6869 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6871 if (!GV || !GV->hasInitializer())
6875 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6878 std::vector<Constant *> NewStructors;
6879 NewStructors.reserve(OldStructorsArray->getNumOperands());
6881 for (
Use &U : OldStructorsArray->operands()) {
6890 Func = UpgradeSinglePointer(Func);
6894 NewStructors.push_back(
6903 if (GlobalArraysToUpgrade.
empty())
6905 assert(UseAddressDisc.has_value());
6907 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6908 GV->setInitializer(NewInit);
6911 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6921 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6925 bool HasObjCFlag =
false, HasClassProperties =
false;
6926 bool HasSwiftVersionFlag =
false;
6927 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6934 if (
Op->getNumOperands() != 3)
6948 if (ID->getString() ==
"Objective-C Image Info Version")
6950 if (ID->getString() ==
"Objective-C Class Properties")
6951 HasClassProperties =
true;
6953 if (ID->getString() ==
"PIC Level") {
6954 if (
auto *Behavior =
6956 uint64_t V = Behavior->getLimitedValue();
6962 if (ID->getString() ==
"PIE Level")
6963 if (
auto *Behavior =
6970 if (ID->getString() ==
"branch-target-enforcement" ||
6971 ID->getString().starts_with(
"sign-return-address")) {
6972 if (
auto *Behavior =
6978 Op->getOperand(1),
Op->getOperand(2)};
6988 if (ID->getString() ==
"Objective-C Image Info Section") {
6991 Value->getString().split(ValueComp,
" ");
6992 if (ValueComp.
size() != 1) {
6993 std::string NewValue;
6994 for (
auto &S : ValueComp)
6995 NewValue += S.str();
7006 if (ID->getString() ==
"Objective-C Garbage Collection") {
7009 assert(Md->getValue() &&
"Expected non-empty metadata");
7010 auto Type = Md->getValue()->getType();
7013 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7014 if ((Val & 0xff) != Val) {
7015 HasSwiftVersionFlag =
true;
7016 SwiftABIVersion = (Val & 0xff00) >> 8;
7017 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7018 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7029 if (ID->getString() ==
"amdgpu_code_object_version") {
7032 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7041 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7070 if (HasObjCFlag && !HasClassProperties) {
7076 if (HasSwiftVersionFlag) {
7080 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7082 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7090 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7094 auto MatchesVersion = [](
const MDNode *
Op) {
7095 return Op->getNumOperands() >= 3 &&
7109 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7110 assert(
Op->getNumOperands() >= 2 &&
7111 "Expected at least 2 operands - name and linkage type");
7123 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7134 auto TrimSpaces = [](
StringRef Section) -> std::string {
7136 Section.split(Components,
',');
7141 for (
auto Component : Components)
7142 OS <<
',' << Component.trim();
7147 for (
auto &GV : M.globals()) {
7148 if (!GV.hasSection())
7153 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7158 GV.setSection(TrimSpaces(Section));
7174struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7175 StrictFPUpgradeVisitor() =
default;
7178 if (!
Call.isStrictFP())
7184 Call.removeFnAttr(Attribute::StrictFP);
7185 Call.addFnAttr(Attribute::NoBuiltin);
7190struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7191 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7192 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7194 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7209 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7210 StrictFPUpgradeVisitor SFPV;
7215 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7216 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7217 for (
auto &Arg :
F.args())
7219 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7221 bool AddingAttrs =
false, RemovingAttrs =
false;
7222 AttrBuilder AttrsToAdd(
F.getContext());
7227 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7228 A.isValid() &&
A.isStringAttribute()) {
7229 F.setSection(
A.getValueAsString());
7231 RemovingAttrs =
true;
7235 A.isValid() &&
A.isStringAttribute()) {
7238 AddingAttrs = RemovingAttrs =
true;
7241 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7242 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7244 RemovingAttrs =
true;
7245 if (
A.getValueAsString() ==
"true") {
7246 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7255 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7258 if (
A.getValueAsBool()) {
7259 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7265 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7266 RemovingAttrs =
true;
7273 bool HandleDenormalMode =
false;
7275 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7278 DenormalFPMath = ParsedMode;
7280 AddingAttrs = RemovingAttrs =
true;
7281 HandleDenormalMode =
true;
7285 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7289 DenormalFPMathF32 = ParsedMode;
7291 AddingAttrs = RemovingAttrs =
true;
7292 HandleDenormalMode =
true;
7296 if (HandleDenormalMode)
7297 AttrsToAdd.addDenormalFPEnvAttr(
7301 F.removeFnAttrs(AttrsToRemove);
7304 F.addFnAttrs(AttrsToAdd);
7310 if (!
F.hasFnAttribute(FnAttrName))
7311 F.addFnAttr(FnAttrName,
Value);
7318 if (!
F.hasFnAttribute(FnAttrName)) {
7320 F.addFnAttr(FnAttrName);
7322 auto A =
F.getFnAttribute(FnAttrName);
7323 if (
"false" ==
A.getValueAsString())
7324 F.removeFnAttr(FnAttrName);
7325 else if (
"true" ==
A.getValueAsString()) {
7326 F.removeFnAttr(FnAttrName);
7327 F.addFnAttr(FnAttrName);
7333 Triple T(M.getTargetTriple());
7334 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7337 uint64_t BTEValue = 0;
7338 uint64_t BPPLRValue = 0;
7339 uint64_t GCSValue = 0;
7340 uint64_t SRAValue = 0;
7341 uint64_t SRAALLValue = 0;
7342 uint64_t SRABKeyValue = 0;
7344 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7348 if (
Op->getNumOperands() != 3)
7357 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7358 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7359 : IDStr ==
"guarded-control-stack" ? &GCSValue
7360 : IDStr ==
"sign-return-address" ? &SRAValue
7361 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7362 : IDStr ==
"sign-return-address-with-bkey"
7368 *ValPtr = CI->getZExtValue();
7374 bool BTE = BTEValue == 1;
7375 bool BPPLR = BPPLRValue == 1;
7376 bool GCS = GCSValue == 1;
7377 bool SRA = SRAValue == 1;
7380 if (SRA && SRAALLValue == 1)
7381 SignTypeValue =
"all";
7384 if (SRA && SRABKeyValue == 1)
7385 SignKeyValue =
"b_key";
7387 for (
Function &
F : M.getFunctionList()) {
7388 if (
F.isDeclaration())
7395 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7396 A.isValid() &&
"none" ==
A.getValueAsString()) {
7397 F.removeFnAttr(
"sign-return-address");
7398 F.removeFnAttr(
"sign-return-address-key");
7414 if (SRAALLValue == 1)
7416 if (SRABKeyValue == 1)
7443 if (
T->getNumOperands() < 1)
7448 if (S->getString().starts_with(
"llvm.vectorizer."))
7454 StringRef OldPrefix =
"llvm.vectorizer.";
7457 if (OldTag ==
"llvm.vectorizer.unroll")
7469 if (
T->getNumOperands() < 1)
7481 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7494 Ops.reserve(
T->getNumOperands());
7495 Ops.push_back(NewTag);
7496 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7497 Ops.push_back(
T->getOperand(
I));
7514 if (
T->isDistinct()) {
7515 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7527 Ops.reserve(
T->getNumOperands());
7538 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7539 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7540 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7543 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7545 auto I =
DL.find(
"-n64-");
7547 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7552 std::string Res =
DL.str();
7555 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7556 Res.append(Res.empty() ?
"G1" :
"-G1");
7564 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7565 Res.append(
"-ni:7:8:9");
7567 if (
DL.ends_with(
"ni:7"))
7569 if (
DL.ends_with(
"ni:7:8"))
7574 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7575 Res.append(
"-p7:160:256:256:32");
7576 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7577 Res.append(
"-p8:128:128:128:48");
7578 constexpr StringRef OldP8(
"-p8:128:128-");
7579 if (
DL.contains(OldP8))
7580 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7581 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7582 Res.append(
"-p9:192:256:256:32");
7586 if (!
DL.contains(
"m:e"))
7587 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7592 if (
T.isSystemZ() && !
DL.empty()) {
7594 if (!
DL.contains(
"-S64"))
7595 return "E-S64" +
DL.drop_front(1).str();
7599 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7602 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7603 if (!
DL.contains(AddrSpaces)) {
7605 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7606 if (R.match(Res, &
Groups))
7612 if (
T.isAArch64()) {
7614 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7615 Res.append(
"-Fn32");
7616 AddPtr32Ptr64AddrSpaces();
7620 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7624 std::string I64 =
"-i64:64";
7625 std::string I128 =
"-i128:128";
7627 size_t Pos = Res.find(I64);
7628 if (Pos !=
size_t(-1))
7629 Res.insert(Pos + I64.size(), I128);
7633 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7634 size_t Pos = Res.find(
"-S128");
7637 Res.insert(Pos,
"-f64:32:64");
7643 AddPtr32Ptr64AddrSpaces();
7651 if (!
T.isOSIAMCU()) {
7652 std::string I128 =
"-i128:128";
7655 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7656 if (R.match(Res, &
Groups))
7664 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7666 auto I =
Ref.find(
"-f80:32-");
7668 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7676 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7679 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7680 B.removeAttribute(
"no-frame-pointer-elim");
7682 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7684 if (FramePointer !=
"all")
7685 FramePointer =
"non-leaf";
7686 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7688 if (!FramePointer.
empty())
7689 B.addAttribute(
"frame-pointer", FramePointer);
7691 A =
B.getAttribute(
"null-pointer-is-valid");
7694 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7695 B.removeAttribute(
"null-pointer-is-valid");
7696 if (NullPointerIsValid)
7697 B.addAttribute(Attribute::NullPointerIsValid);
7700 A =
B.getAttribute(
"uniform-work-group-size");
7704 bool IsTrue = Val ==
"true";
7705 B.removeAttribute(
"uniform-work-group-size");
7707 B.addAttribute(
"uniform-work-group-size");
7718 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static Metadata * upgradeLoopArgument(Metadata *MD)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
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,...
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05MMAIntrinsic(Function *F, StringRef Name)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static cl::opt< bool > DisableAutoUpgradeDebugInfo("disable-auto-upgrade-debug-info", cl::desc("Disable autoupgrade of debug info"))
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static 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.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
StringSwitch & StartsWith(StringLiteral S, T Value)
StringSwitch & Cases(std::initializer_list< StringLiteral > CaseStrings, T Value)
Class to represent struct types.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
unsigned getNumElements() const
Random access to the elements.
Type * getElementType(unsigned N) const
The TimeTraceScope is a helper class to call the begin and end functions of the time trace profiler.
Triple - Helper class for working with autoconf configuration names.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)
bool isVectorTy() const
True if this is an instance of VectorType.
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
LLVM_ABI unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)
LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isPtrOrPtrVectorTy() const
Return true if this is a pointer type or a vector of pointer types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isFPOrFPVectorTy() const
Return true if this is a FP type or a vector of FP.
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
static LLVM_ABI Type * getBFloatTy(LLVMContext &C)
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
A Use represents the edge between a Value definition and its users.
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
LLVM_ABI void print(raw_ostream &O, bool IsForDebug=false) const
Implement operator<< on Value.
LLVM_ABI void setName(const Twine &Name)
Change the name of the value.
LLVM_ABI void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
LLVMContext & getContext() const
All values hold a context through their type.
iterator_range< user_iterator > users()
LLVM_ABI const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
LLVM_ABI void takeName(Value *V)
Transfer the name from V to this value.
Base class of all SIMD vector types.
static VectorType * getInteger(VectorType *VTy)
This static method gets a VectorType with the same number of elements as the input type,...
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
constexpr ScalarTy getFixedValue() const
const ParentTy * getParent() const
self_iterator getIterator()
A raw_ostream that writes to an SmallVector or SmallString.
StringRef str() const
Return a StringRef for the vector contents.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ LOCAL_ADDRESS
Address space for local memory.
@ FLAT_ADDRESS
Address space for flat memory.
@ PRIVATE_ADDRESS
Address space for private memory.
@ PTX_Kernel
Call to a PTX kernel. Passes all arguments in parameter space.
std::optional< ABIType > parseABIType(StringRef S)
Parse the string spelling used by the "float-abi" IR module flag into an ABIType.
LLVM_ABI std::optional< Function * > remangleIntrinsicFunction(Function *F)
LLVM_ABI Function * getOrInsertDeclaration(Module *M, ID id, ArrayRef< Type * > OverloadTys={})
Look up the Function declaration of the intrinsic id in the Module M.
LLVM_ABI ID lookupIntrinsicID(StringRef Name)
This does the actual lookup of an intrinsic ID which matches the given function name.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
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
@ Dynamic
Denotes mode unknown at compile time.
ArrayRef(const T &OneElt) -> ArrayRef< T >
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
@ Default
The result value is uniform if and only if all operands are uniform.
LLVM_ABI bool verifyModule(const Module &M, raw_ostream *OS=nullptr, bool *BrokenDebugInfo=nullptr)
Check a module for errors.
LLVM_ABI void reportFatalUsageError(Error Err)
Report a fatal error that does not indicate a bug in LLVM.
void swap(llvm::BitVector &LHS, llvm::BitVector &RHS)
Implement std::swap in terms of BitVector swap.
This struct is a compact representation of a valid (non-zero power of two) alignment.
Represents the full denormal controls for a function, including the default mode and the f32 specific...
Represent subnormal handling kind for floating point instruction inputs and outputs.
static constexpr DenormalMode getInvalid()
constexpr bool isValid() const
static constexpr DenormalMode getIEEE()
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.