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(
"addqv")) {
1052 if (!
F->getReturnType()->isFPOrFPVectorTy())
1055 auto Args =
F->getFunctionType()->params();
1056 Type *Tys[] = {
F->getReturnType(), Args[1]};
1058 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1062 if (Name.consume_front(
"ld")) {
1064 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1065 if (LdRegex.
match(Name)) {
1071 "Expected 2 arguments for ld* intrinsic.");
1072 Type *PtrTy =
F->getArg(1)->getType();
1075 Intrinsic::aarch64_sve_ld2_sret,
1076 Intrinsic::aarch64_sve_ld3_sret,
1077 Intrinsic::aarch64_sve_ld4_sret,
1080 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1086 if (Name.consume_front(
"tuple.")) {
1088 if (Name.starts_with(
"get")) {
1090 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1092 F->getParent(), Intrinsic::vector_extract, Tys);
1096 if (Name.starts_with(
"set")) {
1098 auto Args =
F->getFunctionType()->params();
1099 Type *Tys[] = {Args[0], Args[2], Args[1]};
1101 F->getParent(), Intrinsic::vector_insert, Tys);
1105 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1106 if (CreateTupleRegex.
match(Name)) {
1108 auto Args =
F->getFunctionType()->params();
1109 Type *Tys[] = {
F->getReturnType(), Args[1]};
1111 F->getParent(), Intrinsic::vector_insert, Tys);
1117 if (Name.starts_with(
"rev.nxv")) {
1120 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1126 if (Name.consume_front(
"sme.")) {
1128 if (Name.consume_front(
"ftmopa.")) {
1133 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1134 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1151 if (Name.consume_front(
"cp.async.bulk.tensor.g2s.")) {
1155 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d)
1157 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d)
1159 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d)
1160 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d)
1161 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d)
1162 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d)
1163 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d)
1164 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d)
1173 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1187 size_t FlagStartIndex =
F->getFunctionType()->getNumParams() - 3;
1188 Type *ArgType =
F->getFunctionType()->getParamType(FlagStartIndex);
1213 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1216 auto [RedOpName, ShapeName] = Name.split(
'.');
1221 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1222 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1223 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1224 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1225 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1226 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1227 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1228 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1234 if (Name.consume_front(
"mapa.shared.cluster"))
1235 if (
F->getReturnType()->getPointerAddressSpace() ==
1237 return Intrinsic::nvvm_mapa_shared_cluster;
1239 if (Name.consume_front(
"cp.async.bulk.")) {
1242 .
Case(
"global.to.shared.cluster",
1243 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1244 .
Case(
"shared.cta.to.cluster",
1245 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1249 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1259 if (!Name.consume_front(
"tcgen05.commit."))
1262 if (Name.consume_front(
"shared."))
1264 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1265 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1268 if (Name.consume_front(
"mc.shared.")) {
1270 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1274 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1275 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1284 if (
F->arg_size() != 2)
1287 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1288 Name.consume_front(
"tcgen05.alloc."))
1290 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1291 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1294 if (Name.consume_front(
"tcgen05.dealloc."))
1296 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1297 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1304 if (Name.consume_front(
"fma.rn."))
1306 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1307 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1308 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1309 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1312 if (Name.consume_front(
"fmax."))
1314 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1315 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1316 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1317 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1318 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1319 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1320 .
Case(
"ftz.nan.xorsign.abs.bf16",
1321 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1322 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1323 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1324 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1325 .
Case(
"ftz.xorsign.abs.bf16x2",
1326 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1327 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1328 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1329 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1330 .
Case(
"nan.xorsign.abs.bf16x2",
1331 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1332 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1333 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1336 if (Name.consume_front(
"fmin."))
1338 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1339 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1340 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1341 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1342 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1343 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1344 .
Case(
"ftz.nan.xorsign.abs.bf16",
1345 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1346 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1347 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1348 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1349 .
Case(
"ftz.xorsign.abs.bf16x2",
1350 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1351 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1352 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1353 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1354 .
Case(
"nan.xorsign.abs.bf16x2",
1355 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1356 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1357 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1360 if (Name.consume_front(
"neg."))
1362 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1363 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1371 if (!Name.consume_front(
"tcgen05.mma."))
1375 if (Name.starts_with(
"ws"))
1378 return F->getIntrinsicID();
1382 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1383 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1384 Name.consume_front(
"param");
1388 if (!Name.consume_front(
"vp."))
1417 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1418 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1425 if (!Name.consume_front(
"vp."))
1445 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1446 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1451 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1463 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1474 if (Name.starts_with(
"to.fp16")) {
1478 FuncTy->getReturnType());
1481 if (Name.starts_with(
"from.fp16")) {
1485 FuncTy->getReturnType());
1497 if (Defaults.empty())
1509 if (
F->arg_size() >= FullDecl->
arg_size())
1514 if (
F->arg_size() < FirstDefault)
1522 bool CanUpgradeDebugIntrinsicsToRecords) {
1523 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1528 if (!Name.consume_front(
"llvm.") || Name.empty())
1534 bool IsArm = Name.consume_front(
"arm.");
1535 if (IsArm || Name.consume_front(
"aarch64.")) {
1541 if (Name.consume_front(
"amdgcn.")) {
1542 if (Name ==
"alignbit") {
1545 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1549 if (Name.consume_front(
"atomic.")) {
1550 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1551 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1560 switch (
F->getIntrinsicID()) {
1564 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1565 if (
F->arg_size() == 7) {
1570 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1571 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1572 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1573 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1574 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1575 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1576 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1577 if (
F->arg_size() == 8) {
1584 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1585 Name.consume_front(
"flat.atomic.")) {
1586 if (Name.starts_with(
"fadd") ||
1588 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1589 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1597 if (Name.starts_with(
"ldexp.")) {
1600 F->getParent(), Intrinsic::ldexp,
1601 {F->getReturnType(), F->getArg(1)->getType()});
1610 if (
F->arg_size() == 1) {
1611 if (Name.consume_front(
"convert.")) {
1625 F->arg_begin()->getType());
1631 if (Name ==
"coro.end" &&
1632 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1633 CoroEndID = Intrinsic::coro_end;
1634 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1635 CoroEndID = Intrinsic::coro_end_async;
1646 if (Name.consume_front(
"dbg.")) {
1648 if (CanUpgradeDebugIntrinsicsToRecords) {
1649 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1650 Name ==
"declare" || Name ==
"label") {
1659 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1662 Intrinsic::dbg_value);
1669 if (Name.consume_front(
"experimental.vector.")) {
1675 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1676 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1677 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1678 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1679 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1681 Intrinsic::vector_partial_reduce_add)
1684 const auto *FT =
F->getFunctionType();
1686 if (ID == Intrinsic::vector_extract ||
1687 ID == Intrinsic::vector_interleave2)
1690 if (ID != Intrinsic::vector_interleave2)
1692 if (ID == Intrinsic::vector_insert ||
1693 ID == Intrinsic::vector_partial_reduce_add)
1701 if (Name.consume_front(
"reduce.")) {
1703 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1704 if (R.match(Name, &
Groups))
1706 .
Case(
"add", Intrinsic::vector_reduce_add)
1707 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1708 .
Case(
"and", Intrinsic::vector_reduce_and)
1709 .
Case(
"or", Intrinsic::vector_reduce_or)
1710 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1711 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1712 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1713 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1714 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1715 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1716 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1721 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1726 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1727 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1732 auto Args =
F->getFunctionType()->params();
1734 {Args[V2 ? 1 : 0]});
1740 if (Name.consume_front(
"splice"))
1744 if (Name.consume_front(
"experimental.stepvector.")) {
1748 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1753 if (Name.starts_with(
"flt.rounds")) {
1756 Intrinsic::get_rounding);
1761 if (Name.starts_with(
"invariant.group.barrier")) {
1763 auto Args =
F->getFunctionType()->params();
1764 Type* ObjectPtr[1] = {Args[0]};
1767 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1772 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1773 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1774 if (IsLifetimeStart || IsLifetimeEnd) {
1775 if (
F->arg_size() == 2) {
1776 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1777 : Intrinsic::lifetime_end;
1782 F->getArg(1)->getType());
1784 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1804 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1805 .StartsWith(
"memmove.", Intrinsic::memmove)
1807 if (
F->arg_size() == 5) {
1811 F->getFunctionType()->params().slice(0, 3);
1817 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1820 const auto *FT =
F->getFunctionType();
1821 Type *ParamTypes[2] = {
1822 FT->getParamType(0),
1826 Intrinsic::memset, ParamTypes);
1832 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1833 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1834 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1835 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1837 if (MaskedID &&
F->arg_size() == 4) {
1839 if (MaskedID == Intrinsic::masked_load ||
1840 MaskedID == Intrinsic::masked_gather) {
1842 F->getParent(), MaskedID,
1843 {F->getReturnType(), F->getArg(0)->getType()});
1847 F->getParent(), MaskedID,
1848 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1854 if (Name.consume_front(
"nvvm.")) {
1856 if (
F->arg_size() == 1) {
1859 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1860 .Case(
"clz.i", Intrinsic::ctlz)
1861 .
Case(
"popc.i", Intrinsic::ctpop)
1865 {F->getReturnType()});
1868 }
else if (
F->arg_size() == 2) {
1871 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1872 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1873 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1874 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1878 {F->getReturnType()});
1884 if (!
F->getReturnType()->getScalarType()->isBFloatTy()) {
1914 F->getParent(), IID,
F->getReturnType(),
1915 F->getFunctionType()->params());
1926 {F->getArg(0)->getType()});
1951 bool Expand =
false;
1952 if (Name.consume_front(
"abs."))
1955 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
1956 else if (Name.consume_front(
"fabs."))
1958 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
1959 else if (Name.consume_front(
"ex2.approx."))
1962 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
1963 else if (Name.consume_front(
"atomic.load."))
1972 else if (Name.consume_front(
"atomic."))
1987 else if (Name.consume_front(
"bitcast."))
1990 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
1991 else if (Name.consume_front(
"rotate."))
1993 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
1994 else if (Name.consume_front(
"ptr.gen.to."))
1997 else if (Name.consume_front(
"ptr."))
2000 else if (Name.consume_front(
"ldg.global."))
2002 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2003 Name.starts_with(
"p."));
2006 .
Case(
"barrier0",
true)
2007 .
Case(
"barrier.n",
true)
2008 .
Case(
"barrier.sync.cnt",
true)
2009 .
Case(
"barrier.sync",
true)
2010 .
Case(
"barrier",
true)
2011 .
Case(
"bar.sync",
true)
2012 .
Case(
"barrier0.popc",
true)
2013 .
Case(
"barrier0.and",
true)
2014 .
Case(
"barrier0.or",
true)
2015 .
Case(
"clz.ll",
true)
2016 .
Case(
"popc.ll",
true)
2018 .
Case(
"swap.lo.hi.b64",
true)
2019 .
Case(
"tanh.approx.f32",
true)
2031 if (Name.starts_with(
"objectsize.")) {
2032 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2033 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2036 Intrinsic::objectsize, Tys);
2043 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2046 F->getParent(), Intrinsic::ptr_annotation,
2047 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2053 if (Name.consume_front(
"riscv.")) {
2056 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2057 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2058 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2059 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2062 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2075 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2076 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2085 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2086 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2087 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2088 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2093 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2102 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2104 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2113 if (Name ==
"stackprotectorcheck") {
2120 if (Name ==
"thread.pointer") {
2122 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2128 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2131 F->getParent(), Intrinsic::var_annotation,
2132 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2135 if (Name.consume_front(
"vector.splice")) {
2136 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2146 if (Name.consume_front(
"wasm.")) {
2149 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2150 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2151 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2156 F->getReturnType());
2160 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2162 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2164 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2183 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2193 std::string
Name =
F->getName().str();
2196 Name,
F->getParent());
2207 if (Result != std::nullopt) {
2223 bool CanUpgradeDebugIntrinsicsToRecords) {
2243 GV->
getName() ==
"llvm.global_dtors")) ||
2258 unsigned N =
Init->getNumOperands();
2259 std::vector<Constant *> NewCtors(
N);
2260 for (
unsigned i = 0; i !=
N; ++i) {
2263 Ctor->getAggregateElement(1),
2277 unsigned NumElts = ResultTy->getNumElements() * 8;
2281 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2291 for (
unsigned l = 0; l != NumElts; l += 16)
2292 for (
unsigned i = 0; i != 16; ++i) {
2293 unsigned Idx = NumElts + i - Shift;
2295 Idx -= NumElts - 16;
2296 Idxs[l + i] = Idx + l;
2299 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2303 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2311 unsigned NumElts = ResultTy->getNumElements() * 8;
2315 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2325 for (
unsigned l = 0; l != NumElts; l += 16)
2326 for (
unsigned i = 0; i != 16; ++i) {
2327 unsigned Idx = i + Shift;
2329 Idx += NumElts - 16;
2330 Idxs[l + i] = Idx + l;
2333 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2337 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2345 Mask = Builder.CreateBitCast(Mask, MaskTy);
2351 for (
unsigned i = 0; i != NumElts; ++i)
2353 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2364 if (
C->isAllOnesValue())
2369 return Builder.CreateSelect(Mask, Op0, Op1);
2376 if (
C->isAllOnesValue())
2380 Mask->getType()->getIntegerBitWidth());
2381 Mask = Builder.CreateBitCast(Mask, MaskTy);
2382 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2383 return Builder.CreateSelect(Mask, Op0, Op1);
2396 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2397 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2402 ShiftVal &= (NumElts - 1);
2411 if (ShiftVal > 16) {
2419 for (
unsigned l = 0; l < NumElts; l += 16) {
2420 for (
unsigned i = 0; i != 16; ++i) {
2421 unsigned Idx = ShiftVal + i;
2422 if (!IsVALIGN && Idx >= 16)
2423 Idx += NumElts - 16;
2424 Indices[l + i] = Idx + l;
2429 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2435 bool ZeroMask,
bool IndexForm) {
2438 unsigned EltWidth = Ty->getScalarSizeInBits();
2439 bool IsFloat = Ty->isFPOrFPVectorTy();
2441 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2442 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2443 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2444 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2445 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2446 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2447 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2448 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2449 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2450 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2451 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2452 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2453 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2454 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2455 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2456 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2457 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2458 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2459 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2460 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2461 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2462 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2463 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2464 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2465 else if (VecWidth == 128 && EltWidth == 16)
2466 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2467 else if (VecWidth == 256 && EltWidth == 16)
2468 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2469 else if (VecWidth == 512 && EltWidth == 16)
2470 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2471 else if (VecWidth == 128 && EltWidth == 8)
2472 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2473 else if (VecWidth == 256 && EltWidth == 8)
2474 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2475 else if (VecWidth == 512 && EltWidth == 8)
2476 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2487 Value *V = Builder.CreateIntrinsic(IID, Args);
2499 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2510 bool IsRotateRight) {
2520 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2521 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2524 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2525 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2570 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2575 bool IsShiftRight,
bool ZeroMask) {
2589 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2590 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2593 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2594 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2609 const Align Alignment =
2611 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2616 if (
C->isAllOnesValue())
2617 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2622 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2628 const Align Alignment =
2637 if (
C->isAllOnesValue())
2638 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2643 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2649 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2650 {Op0, Builder.getInt1(
false)});
2665 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2666 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2667 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2668 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2669 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2672 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2673 LHS = Builder.CreateAnd(
LHS, Mask);
2674 RHS = Builder.CreateAnd(
RHS, Mask);
2691 if (!
C || !
C->isAllOnesValue())
2692 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2697 for (
unsigned i = 0; i != NumElts; ++i)
2699 for (
unsigned i = NumElts; i != 8; ++i)
2700 Indices[i] = NumElts + i % NumElts;
2701 Vec = Builder.CreateShuffleVector(Vec,
2705 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2709 unsigned CC,
bool Signed) {
2717 }
else if (CC == 7) {
2753 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2754 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2756 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2757 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2766 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2772 Name = Name.substr(12);
2777 if (Name.starts_with(
"max.p")) {
2778 if (VecWidth == 128 && EltWidth == 32)
2779 IID = Intrinsic::x86_sse_max_ps;
2780 else if (VecWidth == 128 && EltWidth == 64)
2781 IID = Intrinsic::x86_sse2_max_pd;
2782 else if (VecWidth == 256 && EltWidth == 32)
2783 IID = Intrinsic::x86_avx_max_ps_256;
2784 else if (VecWidth == 256 && EltWidth == 64)
2785 IID = Intrinsic::x86_avx_max_pd_256;
2788 }
else if (Name.starts_with(
"min.p")) {
2789 if (VecWidth == 128 && EltWidth == 32)
2790 IID = Intrinsic::x86_sse_min_ps;
2791 else if (VecWidth == 128 && EltWidth == 64)
2792 IID = Intrinsic::x86_sse2_min_pd;
2793 else if (VecWidth == 256 && EltWidth == 32)
2794 IID = Intrinsic::x86_avx_min_ps_256;
2795 else if (VecWidth == 256 && EltWidth == 64)
2796 IID = Intrinsic::x86_avx_min_pd_256;
2799 }
else if (Name.starts_with(
"pshuf.b.")) {
2800 if (VecWidth == 128)
2801 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2802 else if (VecWidth == 256)
2803 IID = Intrinsic::x86_avx2_pshuf_b;
2804 else if (VecWidth == 512)
2805 IID = Intrinsic::x86_avx512_pshuf_b_512;
2808 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2809 if (VecWidth == 128)
2810 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2811 else if (VecWidth == 256)
2812 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2813 else if (VecWidth == 512)
2814 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2817 }
else if (Name.starts_with(
"pmulh.w.")) {
2818 if (VecWidth == 128)
2819 IID = Intrinsic::x86_sse2_pmulh_w;
2820 else if (VecWidth == 256)
2821 IID = Intrinsic::x86_avx2_pmulh_w;
2822 else if (VecWidth == 512)
2823 IID = Intrinsic::x86_avx512_pmulh_w_512;
2826 }
else if (Name.starts_with(
"pmulhu.w.")) {
2827 if (VecWidth == 128)
2828 IID = Intrinsic::x86_sse2_pmulhu_w;
2829 else if (VecWidth == 256)
2830 IID = Intrinsic::x86_avx2_pmulhu_w;
2831 else if (VecWidth == 512)
2832 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2835 }
else if (Name.starts_with(
"pmaddw.d.")) {
2836 if (VecWidth == 128)
2837 IID = Intrinsic::x86_sse2_pmadd_wd;
2838 else if (VecWidth == 256)
2839 IID = Intrinsic::x86_avx2_pmadd_wd;
2840 else if (VecWidth == 512)
2841 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2844 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2845 if (VecWidth == 128)
2846 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2847 else if (VecWidth == 256)
2848 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2849 else if (VecWidth == 512)
2850 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2853 }
else if (Name.starts_with(
"packsswb.")) {
2854 if (VecWidth == 128)
2855 IID = Intrinsic::x86_sse2_packsswb_128;
2856 else if (VecWidth == 256)
2857 IID = Intrinsic::x86_avx2_packsswb;
2858 else if (VecWidth == 512)
2859 IID = Intrinsic::x86_avx512_packsswb_512;
2862 }
else if (Name.starts_with(
"packssdw.")) {
2863 if (VecWidth == 128)
2864 IID = Intrinsic::x86_sse2_packssdw_128;
2865 else if (VecWidth == 256)
2866 IID = Intrinsic::x86_avx2_packssdw;
2867 else if (VecWidth == 512)
2868 IID = Intrinsic::x86_avx512_packssdw_512;
2871 }
else if (Name.starts_with(
"packuswb.")) {
2872 if (VecWidth == 128)
2873 IID = Intrinsic::x86_sse2_packuswb_128;
2874 else if (VecWidth == 256)
2875 IID = Intrinsic::x86_avx2_packuswb;
2876 else if (VecWidth == 512)
2877 IID = Intrinsic::x86_avx512_packuswb_512;
2880 }
else if (Name.starts_with(
"packusdw.")) {
2881 if (VecWidth == 128)
2882 IID = Intrinsic::x86_sse41_packusdw;
2883 else if (VecWidth == 256)
2884 IID = Intrinsic::x86_avx2_packusdw;
2885 else if (VecWidth == 512)
2886 IID = Intrinsic::x86_avx512_packusdw_512;
2889 }
else if (Name.starts_with(
"vpermilvar.")) {
2890 if (VecWidth == 128 && EltWidth == 32)
2891 IID = Intrinsic::x86_avx_vpermilvar_ps;
2892 else if (VecWidth == 128 && EltWidth == 64)
2893 IID = Intrinsic::x86_avx_vpermilvar_pd;
2894 else if (VecWidth == 256 && EltWidth == 32)
2895 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2896 else if (VecWidth == 256 && EltWidth == 64)
2897 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2898 else if (VecWidth == 512 && EltWidth == 32)
2899 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2900 else if (VecWidth == 512 && EltWidth == 64)
2901 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2904 }
else if (Name ==
"cvtpd2dq.256") {
2905 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2906 }
else if (Name ==
"cvtpd2ps.256") {
2907 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2908 }
else if (Name ==
"cvttpd2dq.256") {
2909 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2910 }
else if (Name ==
"cvttps2dq.128") {
2911 IID = Intrinsic::x86_sse2_cvttps2dq;
2912 }
else if (Name ==
"cvttps2dq.256") {
2913 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2914 }
else if (Name.starts_with(
"permvar.")) {
2916 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2917 IID = Intrinsic::x86_avx2_permps;
2918 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2919 IID = Intrinsic::x86_avx2_permd;
2920 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2921 IID = Intrinsic::x86_avx512_permvar_df_256;
2922 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2923 IID = Intrinsic::x86_avx512_permvar_di_256;
2924 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2925 IID = Intrinsic::x86_avx512_permvar_sf_512;
2926 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2927 IID = Intrinsic::x86_avx512_permvar_si_512;
2928 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2929 IID = Intrinsic::x86_avx512_permvar_df_512;
2930 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2931 IID = Intrinsic::x86_avx512_permvar_di_512;
2932 else if (VecWidth == 128 && EltWidth == 16)
2933 IID = Intrinsic::x86_avx512_permvar_hi_128;
2934 else if (VecWidth == 256 && EltWidth == 16)
2935 IID = Intrinsic::x86_avx512_permvar_hi_256;
2936 else if (VecWidth == 512 && EltWidth == 16)
2937 IID = Intrinsic::x86_avx512_permvar_hi_512;
2938 else if (VecWidth == 128 && EltWidth == 8)
2939 IID = Intrinsic::x86_avx512_permvar_qi_128;
2940 else if (VecWidth == 256 && EltWidth == 8)
2941 IID = Intrinsic::x86_avx512_permvar_qi_256;
2942 else if (VecWidth == 512 && EltWidth == 8)
2943 IID = Intrinsic::x86_avx512_permvar_qi_512;
2946 }
else if (Name.starts_with(
"dbpsadbw.")) {
2947 if (VecWidth == 128)
2948 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2949 else if (VecWidth == 256)
2950 IID = Intrinsic::x86_avx512_dbpsadbw_256;
2951 else if (VecWidth == 512)
2952 IID = Intrinsic::x86_avx512_dbpsadbw_512;
2955 }
else if (Name.starts_with(
"pmultishift.qb.")) {
2956 if (VecWidth == 128)
2957 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
2958 else if (VecWidth == 256)
2959 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
2960 else if (VecWidth == 512)
2961 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
2964 }
else if (Name.starts_with(
"conflict.")) {
2965 if (Name[9] ==
'd' && VecWidth == 128)
2966 IID = Intrinsic::x86_avx512_conflict_d_128;
2967 else if (Name[9] ==
'd' && VecWidth == 256)
2968 IID = Intrinsic::x86_avx512_conflict_d_256;
2969 else if (Name[9] ==
'd' && VecWidth == 512)
2970 IID = Intrinsic::x86_avx512_conflict_d_512;
2971 else if (Name[9] ==
'q' && VecWidth == 128)
2972 IID = Intrinsic::x86_avx512_conflict_q_128;
2973 else if (Name[9] ==
'q' && VecWidth == 256)
2974 IID = Intrinsic::x86_avx512_conflict_q_256;
2975 else if (Name[9] ==
'q' && VecWidth == 512)
2976 IID = Intrinsic::x86_avx512_conflict_q_512;
2979 }
else if (Name.starts_with(
"pavg.")) {
2980 if (Name[5] ==
'b' && VecWidth == 128)
2981 IID = Intrinsic::x86_sse2_pavg_b;
2982 else if (Name[5] ==
'b' && VecWidth == 256)
2983 IID = Intrinsic::x86_avx2_pavg_b;
2984 else if (Name[5] ==
'b' && VecWidth == 512)
2985 IID = Intrinsic::x86_avx512_pavg_b_512;
2986 else if (Name[5] ==
'w' && VecWidth == 128)
2987 IID = Intrinsic::x86_sse2_pavg_w;
2988 else if (Name[5] ==
'w' && VecWidth == 256)
2989 IID = Intrinsic::x86_avx2_pavg_w;
2990 else if (Name[5] ==
'w' && VecWidth == 512)
2991 IID = Intrinsic::x86_avx512_pavg_w_512;
3000 Rep = Builder.CreateIntrinsic(IID, Args);
3011 if (AsmStr->find(
"mov\tfp") == 0 &&
3012 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3013 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3014 AsmStr->replace(Pos, 1,
";");
3020 Value *Rep =
nullptr;
3022 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3024 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3025 {Arg, Builder.getTrue()},
3027 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3028 Type *Ty = (Name ==
"abs.bf16")
3032 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3033 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3034 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3035 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3036 : Intrinsic::nvvm_fabs;
3037 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3038 }
else if (Name.consume_front(
"ex2.approx.")) {
3040 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3041 : Intrinsic::nvvm_ex2_approx;
3042 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3043 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3044 Name.starts_with(
"atomic.load.add.f64.p")) {
3047 Rep = Builder.CreateAtomicRMW(
3053 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3054 Name.starts_with(
"atomic.load.dec.32.p")) {
3059 Rep = Builder.CreateAtomicRMW(
3063 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3069 Op.contains(
".cta.") ?
"block" :
"");
3070 if (
Op.starts_with(
"cas.")) {
3072 Value *Pair = Builder.CreateAtomicCmpXchg(
3075 Rep = Builder.CreateExtractValue(Pair, 0);
3093 "unexpected nvvm scoped atomic intrinsic");
3094 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3097 }
else if (Name ==
"clz.ll") {
3100 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3101 {Arg, Builder.getFalse()},
3103 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3104 }
else if (Name ==
"popc.ll") {
3108 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3109 Arg,
nullptr,
"ctpop");
3110 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3111 }
else if (Name ==
"h2f") {
3113 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3114 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3115 }
else if (Name.consume_front(
"bitcast.") &&
3116 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3119 }
else if (Name ==
"rotate.b32") {
3122 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3123 {Arg, Arg, ShiftAmt});
3124 }
else if (Name ==
"rotate.b64") {
3128 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3129 {Arg, Arg, ZExtShiftAmt});
3130 }
else if (Name ==
"rotate.right.b64") {
3134 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3135 {Arg, Arg, ZExtShiftAmt});
3136 }
else if (Name ==
"swap.lo.hi.b64") {
3139 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3140 {Arg, Arg, Builder.getInt64(32)});
3141 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3144 Name.starts_with(
".to.gen"))) {
3146 }
else if (Name.consume_front(
"ldg.global")) {
3150 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3153 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3155 }
else if (Name ==
"tanh.approx.f32") {
3159 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3161 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3163 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3164 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3166 }
else if (Name ==
"barrier") {
3167 Rep = Builder.CreateIntrinsic(
3168 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3170 }
else if (Name ==
"barrier.sync") {
3171 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3173 }
else if (Name ==
"barrier.sync.cnt") {
3174 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3176 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3177 Name ==
"barrier0.or") {
3179 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3183 .
Case(
"barrier0.popc",
3184 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3185 .
Case(
"barrier0.and",
3186 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3187 .
Case(
"barrier0.or",
3188 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3189 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3190 Rep = Builder.CreateZExt(Bar, CI->
getType());
3194 !
F->getReturnType()->getScalarType()->isBFloatTy()) {
3204 ? Builder.CreateBitCast(Arg, NewType)
3207 Rep = Builder.CreateCall(NewFn, Args);
3208 if (
F->getReturnType()->isIntegerTy())
3209 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3219 Value *Rep =
nullptr;
3221 if (Name.starts_with(
"sse4a.movnt.")) {
3233 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3236 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3237 }
else if (Name.starts_with(
"avx.movnt.") ||
3238 Name.starts_with(
"avx512.storent.")) {
3250 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3251 }
else if (Name ==
"sse2.storel.dq") {
3256 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3257 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3258 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3259 }
else if (Name.starts_with(
"sse.storeu.") ||
3260 Name.starts_with(
"sse2.storeu.") ||
3261 Name.starts_with(
"avx.storeu.")) {
3264 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3265 }
else if (Name ==
"avx512.mask.store.ss") {
3269 }
else if (Name.starts_with(
"avx512.mask.store")) {
3271 bool Aligned = Name[17] !=
'u';
3274 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3277 bool CmpEq = Name[9] ==
'e';
3280 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3281 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3288 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3289 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3291 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3292 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3293 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3294 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3295 Name.starts_with(
"sse2.sqrt.p") ||
3296 Name.starts_with(
"sse.sqrt.p")) {
3297 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3298 {CI->getArgOperand(0)});
3299 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3303 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3304 : Intrinsic::x86_avx512_sqrt_pd_512;
3307 Rep = Builder.CreateIntrinsic(IID, Args);
3309 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3310 {CI->getArgOperand(0)});
3314 }
else if (Name.starts_with(
"avx512.ptestm") ||
3315 Name.starts_with(
"avx512.ptestnm")) {
3319 Rep = Builder.CreateAnd(Op0, Op1);
3325 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3327 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3330 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3333 }
else if (Name.starts_with(
"avx512.kunpck")) {
3338 for (
unsigned i = 0; i != NumElts; ++i)
3347 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3348 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3349 }
else if (Name ==
"avx512.kand.w") {
3352 Rep = Builder.CreateAnd(
LHS,
RHS);
3353 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3354 }
else if (Name ==
"avx512.kandn.w") {
3357 LHS = Builder.CreateNot(
LHS);
3358 Rep = Builder.CreateAnd(
LHS,
RHS);
3359 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3360 }
else if (Name ==
"avx512.kor.w") {
3363 Rep = Builder.CreateOr(
LHS,
RHS);
3364 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3365 }
else if (Name ==
"avx512.kxor.w") {
3368 Rep = Builder.CreateXor(
LHS,
RHS);
3369 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3370 }
else if (Name ==
"avx512.kxnor.w") {
3373 LHS = Builder.CreateNot(
LHS);
3374 Rep = Builder.CreateXor(
LHS,
RHS);
3375 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3376 }
else if (Name ==
"avx512.knot.w") {
3378 Rep = Builder.CreateNot(Rep);
3379 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3380 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3383 Rep = Builder.CreateOr(
LHS,
RHS);
3384 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3386 if (Name[14] ==
'c')
3390 Rep = Builder.CreateICmpEQ(Rep,
C);
3391 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3392 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3393 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3394 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3395 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3398 ConstantInt::get(I32Ty, 0));
3400 ConstantInt::get(I32Ty, 0));
3402 if (Name.contains(
".add."))
3403 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3404 else if (Name.contains(
".sub."))
3405 EltOp = Builder.CreateFSub(Elt0, Elt1);
3406 else if (Name.contains(
".mul."))
3407 EltOp = Builder.CreateFMul(Elt0, Elt1);
3409 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3410 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3411 ConstantInt::get(I32Ty, 0));
3412 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3414 bool CmpEq = Name[16] ==
'e';
3416 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3418 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3425 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3428 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3431 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3438 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3440 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3441 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3443 if (VecWidth == 128 && EltWidth == 32)
3444 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3445 else if (VecWidth == 256 && EltWidth == 32)
3446 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3447 else if (VecWidth == 512 && EltWidth == 32)
3448 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3449 else if (VecWidth == 128 && EltWidth == 64)
3450 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3451 else if (VecWidth == 256 && EltWidth == 64)
3452 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3453 else if (VecWidth == 512 && EltWidth == 64)
3454 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3461 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3464 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3465 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3467 if (VecWidth == 128 && EltWidth == 32)
3468 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3469 else if (VecWidth == 256 && EltWidth == 32)
3470 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3471 else if (VecWidth == 512 && EltWidth == 32)
3472 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3473 else if (VecWidth == 128 && EltWidth == 64)
3474 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3475 else if (VecWidth == 256 && EltWidth == 64)
3476 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3477 else if (VecWidth == 512 && EltWidth == 64)
3478 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3483 if (VecWidth == 512)
3485 Args.push_back(Mask);
3487 Rep = Builder.CreateIntrinsic(IID, Args);
3488 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3492 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3495 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3496 Name.starts_with(
"avx512.cvtw2mask.") ||
3497 Name.starts_with(
"avx512.cvtd2mask.") ||
3498 Name.starts_with(
"avx512.cvtq2mask.")) {
3503 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3504 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3505 Name.starts_with(
"avx512.mask.pabs")) {
3507 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3508 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3509 Name.starts_with(
"avx512.mask.pmaxs")) {
3511 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3512 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3513 Name.starts_with(
"avx512.mask.pmaxu")) {
3515 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3516 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3517 Name.starts_with(
"avx512.mask.pmins")) {
3519 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3520 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3521 Name.starts_with(
"avx512.mask.pminu")) {
3523 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3524 Name ==
"avx512.pmulu.dq.512" ||
3525 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3527 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3528 Name ==
"avx512.pmul.dq.512" ||
3529 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3531 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3532 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3537 }
else if (Name ==
"avx512.cvtusi2sd") {
3542 }
else if (Name ==
"sse2.cvtss2sd") {
3544 Rep = Builder.CreateFPExt(
3547 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3548 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3549 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3550 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3551 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3552 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3553 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3554 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3555 Name ==
"avx512.mask.cvtqq2ps.256" ||
3556 Name ==
"avx512.mask.cvtqq2ps.512" ||
3557 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3558 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3559 Name ==
"avx.cvt.ps2.pd.256" ||
3560 Name ==
"avx512.mask.cvtps2pd.128" ||
3561 Name ==
"avx512.mask.cvtps2pd.256") {
3566 unsigned NumDstElts = DstTy->getNumElements();
3567 if (NumDstElts < SrcTy->getNumElements()) {
3568 assert(NumDstElts == 2 &&
"Unexpected vector size");
3569 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3572 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3573 bool IsUnsigned = Name.contains(
"cvtu");
3575 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3579 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3580 : Intrinsic::x86_avx512_sitofp_round;
3581 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3584 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3585 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3591 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3592 Name.starts_with(
"vcvtph2ps.")) {
3596 unsigned NumDstElts = DstTy->getNumElements();
3597 if (NumDstElts != SrcTy->getNumElements()) {
3598 assert(NumDstElts == 4 &&
"Unexpected vector size");
3599 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3601 Rep = Builder.CreateBitCast(
3603 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3607 }
else if (Name.starts_with(
"avx512.mask.load")) {
3609 bool Aligned = Name[16] !=
'u';
3612 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3616 ResultTy->getNumElements());
3617 Rep = Builder.CreateIntrinsic(
3618 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3620 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3626 Rep = Builder.CreateIntrinsic(
3627 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3629 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3630 Name.starts_with(
"avx512.mask.expand.")) {
3634 ResultTy->getNumElements());
3636 bool IsCompress = Name[12] ==
'c';
3637 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3638 : Intrinsic::x86_avx512_mask_expand;
3639 Rep = Builder.CreateIntrinsic(
3641 }
else if (Name.starts_with(
"xop.vpcom")) {
3643 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3644 Name.ends_with(
"uq"))
3646 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3647 Name.ends_with(
"d") || Name.ends_with(
"q"))
3656 Name = Name.substr(9);
3657 if (Name.starts_with(
"lt"))
3659 else if (Name.starts_with(
"le"))
3661 else if (Name.starts_with(
"gt"))
3663 else if (Name.starts_with(
"ge"))
3665 else if (Name.starts_with(
"eq"))
3667 else if (Name.starts_with(
"ne"))
3669 else if (Name.starts_with(
"false"))
3671 else if (Name.starts_with(
"true"))
3678 }
else if (Name.starts_with(
"xop.vpcmov")) {
3680 Value *NotSel = Builder.CreateNot(Sel);
3683 Rep = Builder.CreateOr(Sel0, Sel1);
3684 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3685 Name.starts_with(
"avx512.mask.prol")) {
3687 }
else if (Name.starts_with(
"avx512.pror") ||
3688 Name.starts_with(
"avx512.mask.pror")) {
3690 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3691 Name.starts_with(
"avx512.mask.vpshld") ||
3692 Name.starts_with(
"avx512.maskz.vpshld")) {
3693 bool ZeroMask = Name[11] ==
'z';
3695 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3696 Name.starts_with(
"avx512.mask.vpshrd") ||
3697 Name.starts_with(
"avx512.maskz.vpshrd")) {
3698 bool ZeroMask = Name[11] ==
'z';
3700 }
else if (Name ==
"sse42.crc32.64.8") {
3703 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3705 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3706 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3707 Name.starts_with(
"avx512.vbroadcast.s")) {
3710 Type *EltTy = VecTy->getElementType();
3711 unsigned EltNum = VecTy->getNumElements();
3715 for (
unsigned I = 0;
I < EltNum; ++
I)
3716 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3717 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3718 Name.starts_with(
"sse41.pmovzx") ||
3719 Name.starts_with(
"avx2.pmovsx") ||
3720 Name.starts_with(
"avx2.pmovzx") ||
3721 Name.starts_with(
"avx512.mask.pmovsx") ||
3722 Name.starts_with(
"avx512.mask.pmovzx")) {
3724 unsigned NumDstElts = DstTy->getNumElements();
3728 for (
unsigned i = 0; i != NumDstElts; ++i)
3733 bool DoSext = Name.contains(
"pmovsx");
3735 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3740 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3741 Name ==
"avx512.mask.pmov.qd.512" ||
3742 Name ==
"avx512.mask.pmov.wb.256" ||
3743 Name ==
"avx512.mask.pmov.wb.512") {
3748 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3749 Name ==
"avx2.vbroadcasti128") {
3755 if (NumSrcElts == 2)
3758 Rep = Builder.CreateShuffleVector(
Load,
3760 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3761 Name.starts_with(
"avx512.mask.shuf.f")) {
3766 unsigned ControlBitsMask = NumLanes - 1;
3767 unsigned NumControlBits = NumLanes / 2;
3770 for (
unsigned l = 0; l != NumLanes; ++l) {
3771 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3773 if (l >= NumLanes / 2)
3774 LaneMask += NumLanes;
3775 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3776 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3782 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3783 Name.starts_with(
"avx512.mask.broadcasti")) {
3786 unsigned NumDstElts =
3790 for (
unsigned i = 0; i != NumDstElts; ++i)
3791 ShuffleMask[i] = i % NumSrcElts;
3797 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3798 Name.starts_with(
"avx2.vbroadcast") ||
3799 Name.starts_with(
"avx512.pbroadcast") ||
3800 Name.starts_with(
"avx512.mask.broadcast.s")) {
3807 Rep = Builder.CreateShuffleVector(
Op, M);
3812 }
else if (Name.starts_with(
"sse2.padds.") ||
3813 Name.starts_with(
"avx2.padds.") ||
3814 Name.starts_with(
"avx512.padds.") ||
3815 Name.starts_with(
"avx512.mask.padds.")) {
3817 }
else if (Name.starts_with(
"sse2.psubs.") ||
3818 Name.starts_with(
"avx2.psubs.") ||
3819 Name.starts_with(
"avx512.psubs.") ||
3820 Name.starts_with(
"avx512.mask.psubs.")) {
3822 }
else if (Name.starts_with(
"sse2.paddus.") ||
3823 Name.starts_with(
"avx2.paddus.") ||
3824 Name.starts_with(
"avx512.mask.paddus.")) {
3826 }
else if (Name.starts_with(
"sse2.psubus.") ||
3827 Name.starts_with(
"avx2.psubus.") ||
3828 Name.starts_with(
"avx512.mask.psubus.")) {
3830 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3835 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3839 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3844 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3849 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3850 Name ==
"avx512.psll.dq.512") {
3854 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3855 Name ==
"avx512.psrl.dq.512") {
3859 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3860 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3861 Name.starts_with(
"avx2.pblendd.")) {
3866 unsigned NumElts = VecTy->getNumElements();
3869 for (
unsigned i = 0; i != NumElts; ++i)
3870 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
3872 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3873 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3874 Name ==
"avx2.vinserti128" ||
3875 Name.starts_with(
"avx512.mask.insert")) {
3879 unsigned DstNumElts =
3881 unsigned SrcNumElts =
3883 unsigned Scale = DstNumElts / SrcNumElts;
3890 for (
unsigned i = 0; i != SrcNumElts; ++i)
3892 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3893 Idxs[i] = SrcNumElts;
3894 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3908 for (
unsigned i = 0; i != DstNumElts; ++i)
3911 for (
unsigned i = 0; i != SrcNumElts; ++i)
3912 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
3913 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3919 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3920 Name ==
"avx2.vextracti128" ||
3921 Name.starts_with(
"avx512.mask.vextract")) {
3924 unsigned DstNumElts =
3926 unsigned SrcNumElts =
3928 unsigned Scale = SrcNumElts / DstNumElts;
3935 for (
unsigned i = 0; i != DstNumElts; ++i) {
3936 Idxs[i] = i + (
Imm * DstNumElts);
3938 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3944 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3945 Name.starts_with(
"avx512.mask.perm.di.")) {
3949 unsigned NumElts = VecTy->getNumElements();
3952 for (
unsigned i = 0; i != NumElts; ++i)
3953 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
3955 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3960 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
3972 unsigned HalfSize = NumElts / 2;
3984 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
3985 for (
unsigned i = 0; i < HalfSize; ++i)
3986 ShuffleMask[i] = StartIndex + i;
3989 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
3990 for (
unsigned i = 0; i < HalfSize; ++i)
3991 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
3993 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
3995 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
3996 Name.starts_with(
"avx512.mask.vpermil.p") ||
3997 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4001 unsigned NumElts = VecTy->getNumElements();
4003 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4004 unsigned IdxMask = ((1 << IdxSize) - 1);
4010 for (
unsigned i = 0; i != NumElts; ++i)
4011 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4013 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4018 }
else if (Name ==
"sse2.pshufl.w" ||
4019 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4024 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4028 for (
unsigned l = 0; l != NumElts; l += 8) {
4029 for (
unsigned i = 0; i != 4; ++i)
4030 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4031 for (
unsigned i = 4; i != 8; ++i)
4032 Idxs[i + l] = i + l;
4035 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4040 }
else if (Name ==
"sse2.pshufh.w" ||
4041 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4046 if (Name ==
"sse2.pshufh.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] = i + l;
4053 for (
unsigned i = 0; i != 4; ++i)
4054 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4057 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4062 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4069 unsigned HalfLaneElts = NumLaneElts / 2;
4072 for (
unsigned i = 0; i != NumElts; ++i) {
4074 Idxs[i] = i - (i % NumLaneElts);
4076 if ((i % NumLaneElts) >= HalfLaneElts)
4080 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4083 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4087 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4088 Name.starts_with(
"avx512.mask.movshdup") ||
4089 Name.starts_with(
"avx512.mask.movsldup")) {
4095 if (Name.starts_with(
"avx512.mask.movshdup."))
4099 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4100 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4101 Idxs[i + l + 0] = i + l +
Offset;
4102 Idxs[i + l + 1] = i + l +
Offset;
4105 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4109 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4110 Name.starts_with(
"avx512.mask.unpckl.")) {
4117 for (
int l = 0; l != NumElts; l += NumLaneElts)
4118 for (
int i = 0; i != NumLaneElts; ++i)
4119 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4121 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4125 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4126 Name.starts_with(
"avx512.mask.unpckh.")) {
4133 for (
int l = 0; l != NumElts; l += NumLaneElts)
4134 for (
int i = 0; i != NumLaneElts; ++i)
4135 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4137 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4141 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4142 Name.starts_with(
"avx512.mask.pand.")) {
4145 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4147 Rep = Builder.CreateBitCast(Rep, FTy);
4150 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4151 Name.starts_with(
"avx512.mask.pandn.")) {
4154 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4155 Rep = Builder.CreateAnd(Rep,
4157 Rep = Builder.CreateBitCast(Rep, FTy);
4160 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4161 Name.starts_with(
"avx512.mask.por.")) {
4164 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4166 Rep = Builder.CreateBitCast(Rep, FTy);
4169 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4170 Name.starts_with(
"avx512.mask.pxor.")) {
4173 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4175 Rep = Builder.CreateBitCast(Rep, FTy);
4178 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4182 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4186 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4190 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4191 if (Name.ends_with(
".512")) {
4193 if (Name[17] ==
's')
4194 IID = Intrinsic::x86_avx512_add_ps_512;
4196 IID = Intrinsic::x86_avx512_add_pd_512;
4198 Rep = Builder.CreateIntrinsic(
4206 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4207 if (Name.ends_with(
".512")) {
4209 if (Name[17] ==
's')
4210 IID = Intrinsic::x86_avx512_div_ps_512;
4212 IID = Intrinsic::x86_avx512_div_pd_512;
4214 Rep = Builder.CreateIntrinsic(
4222 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4223 if (Name.ends_with(
".512")) {
4225 if (Name[17] ==
's')
4226 IID = Intrinsic::x86_avx512_mul_ps_512;
4228 IID = Intrinsic::x86_avx512_mul_pd_512;
4230 Rep = Builder.CreateIntrinsic(
4238 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4239 if (Name.ends_with(
".512")) {
4241 if (Name[17] ==
's')
4242 IID = Intrinsic::x86_avx512_sub_ps_512;
4244 IID = Intrinsic::x86_avx512_sub_pd_512;
4246 Rep = Builder.CreateIntrinsic(
4254 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4255 Name.starts_with(
"avx512.mask.min.p")) &&
4256 Name.drop_front(18) ==
".512") {
4257 bool IsDouble = Name[17] ==
'd';
4258 bool IsMin = Name[13] ==
'i';
4260 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4261 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4264 Rep = Builder.CreateIntrinsic(
4269 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4271 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4272 {CI->getArgOperand(0), Builder.getInt1(false)});
4275 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4276 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4277 bool IsVariable = Name[16] ==
'v';
4278 char Size = Name[16] ==
'.' ? Name[17]
4279 : Name[17] ==
'.' ? Name[18]
4280 : Name[18] ==
'.' ? Name[19]
4284 if (IsVariable && Name[17] !=
'.') {
4285 if (
Size ==
'd' && Name[17] ==
'2')
4286 IID = Intrinsic::x86_avx2_psllv_q;
4287 else if (
Size ==
'd' && Name[17] ==
'4')
4288 IID = Intrinsic::x86_avx2_psllv_q_256;
4289 else if (
Size ==
's' && Name[17] ==
'4')
4290 IID = Intrinsic::x86_avx2_psllv_d;
4291 else if (
Size ==
's' && Name[17] ==
'8')
4292 IID = Intrinsic::x86_avx2_psllv_d_256;
4293 else if (
Size ==
'h' && Name[17] ==
'8')
4294 IID = Intrinsic::x86_avx512_psllv_w_128;
4295 else if (
Size ==
'h' && Name[17] ==
'1')
4296 IID = Intrinsic::x86_avx512_psllv_w_256;
4297 else if (Name[17] ==
'3' && Name[18] ==
'2')
4298 IID = Intrinsic::x86_avx512_psllv_w_512;
4301 }
else if (Name.ends_with(
".128")) {
4303 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4304 : Intrinsic::x86_sse2_psll_d;
4305 else if (
Size ==
'q')
4306 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4307 : Intrinsic::x86_sse2_psll_q;
4308 else if (
Size ==
'w')
4309 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4310 : Intrinsic::x86_sse2_psll_w;
4313 }
else if (Name.ends_with(
".256")) {
4315 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4316 : Intrinsic::x86_avx2_psll_d;
4317 else if (
Size ==
'q')
4318 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4319 : Intrinsic::x86_avx2_psll_q;
4320 else if (
Size ==
'w')
4321 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4322 : Intrinsic::x86_avx2_psll_w;
4327 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4328 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4329 : Intrinsic::x86_avx512_psll_d_512;
4330 else if (
Size ==
'q')
4331 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4332 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4333 : Intrinsic::x86_avx512_psll_q_512;
4334 else if (
Size ==
'w')
4335 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4336 : Intrinsic::x86_avx512_psll_w_512;
4342 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4343 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4344 bool IsVariable = Name[16] ==
'v';
4345 char Size = Name[16] ==
'.' ? Name[17]
4346 : Name[17] ==
'.' ? Name[18]
4347 : Name[18] ==
'.' ? Name[19]
4351 if (IsVariable && Name[17] !=
'.') {
4352 if (
Size ==
'd' && Name[17] ==
'2')
4353 IID = Intrinsic::x86_avx2_psrlv_q;
4354 else if (
Size ==
'd' && Name[17] ==
'4')
4355 IID = Intrinsic::x86_avx2_psrlv_q_256;
4356 else if (
Size ==
's' && Name[17] ==
'4')
4357 IID = Intrinsic::x86_avx2_psrlv_d;
4358 else if (
Size ==
's' && Name[17] ==
'8')
4359 IID = Intrinsic::x86_avx2_psrlv_d_256;
4360 else if (
Size ==
'h' && Name[17] ==
'8')
4361 IID = Intrinsic::x86_avx512_psrlv_w_128;
4362 else if (
Size ==
'h' && Name[17] ==
'1')
4363 IID = Intrinsic::x86_avx512_psrlv_w_256;
4364 else if (Name[17] ==
'3' && Name[18] ==
'2')
4365 IID = Intrinsic::x86_avx512_psrlv_w_512;
4368 }
else if (Name.ends_with(
".128")) {
4370 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4371 : Intrinsic::x86_sse2_psrl_d;
4372 else if (
Size ==
'q')
4373 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4374 : Intrinsic::x86_sse2_psrl_q;
4375 else if (
Size ==
'w')
4376 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4377 : Intrinsic::x86_sse2_psrl_w;
4380 }
else if (Name.ends_with(
".256")) {
4382 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4383 : Intrinsic::x86_avx2_psrl_d;
4384 else if (
Size ==
'q')
4385 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4386 : Intrinsic::x86_avx2_psrl_q;
4387 else if (
Size ==
'w')
4388 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4389 : Intrinsic::x86_avx2_psrl_w;
4394 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4395 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4396 : Intrinsic::x86_avx512_psrl_d_512;
4397 else if (
Size ==
'q')
4398 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4399 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4400 : Intrinsic::x86_avx512_psrl_q_512;
4401 else if (
Size ==
'w')
4402 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4403 : Intrinsic::x86_avx512_psrl_w_512;
4409 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4410 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4411 bool IsVariable = Name[16] ==
'v';
4412 char Size = Name[16] ==
'.' ? Name[17]
4413 : Name[17] ==
'.' ? Name[18]
4414 : Name[18] ==
'.' ? Name[19]
4418 if (IsVariable && Name[17] !=
'.') {
4419 if (
Size ==
's' && Name[17] ==
'4')
4420 IID = Intrinsic::x86_avx2_psrav_d;
4421 else if (
Size ==
's' && Name[17] ==
'8')
4422 IID = Intrinsic::x86_avx2_psrav_d_256;
4423 else if (
Size ==
'h' && Name[17] ==
'8')
4424 IID = Intrinsic::x86_avx512_psrav_w_128;
4425 else if (
Size ==
'h' && Name[17] ==
'1')
4426 IID = Intrinsic::x86_avx512_psrav_w_256;
4427 else if (Name[17] ==
'3' && Name[18] ==
'2')
4428 IID = Intrinsic::x86_avx512_psrav_w_512;
4431 }
else if (Name.ends_with(
".128")) {
4433 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4434 : Intrinsic::x86_sse2_psra_d;
4435 else if (
Size ==
'q')
4436 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4437 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4438 : Intrinsic::x86_avx512_psra_q_128;
4439 else if (
Size ==
'w')
4440 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4441 : Intrinsic::x86_sse2_psra_w;
4444 }
else if (Name.ends_with(
".256")) {
4446 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4447 : Intrinsic::x86_avx2_psra_d;
4448 else if (
Size ==
'q')
4449 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4450 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4451 : Intrinsic::x86_avx512_psra_q_256;
4452 else if (
Size ==
'w')
4453 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4454 : Intrinsic::x86_avx2_psra_w;
4459 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4460 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4461 : Intrinsic::x86_avx512_psra_d_512;
4462 else if (
Size ==
'q')
4463 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4464 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4465 : Intrinsic::x86_avx512_psra_q_512;
4466 else if (
Size ==
'w')
4467 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4468 : Intrinsic::x86_avx512_psra_w_512;
4474 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4476 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4478 }
else if (Name.ends_with(
".movntdqa")) {
4482 LoadInst *LI = Builder.CreateAlignedLoad(
4487 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4488 Name.starts_with(
"fma.vfmsub.") ||
4489 Name.starts_with(
"fma.vfnmadd.") ||
4490 Name.starts_with(
"fma.vfnmsub.")) {
4491 bool NegMul = Name[6] ==
'n';
4492 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4493 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4504 if (NegMul && !IsScalar)
4505 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4506 if (NegMul && IsScalar)
4507 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4509 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4511 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4515 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4523 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4527 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4528 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4529 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4530 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4531 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4532 bool IsMask3 = Name[11] ==
'3';
4533 bool IsMaskZ = Name[11] ==
'z';
4535 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4536 bool NegMul = Name[2] ==
'n';
4537 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4543 if (NegMul && (IsMask3 || IsMaskZ))
4544 A = Builder.CreateFNeg(
A);
4545 if (NegMul && !(IsMask3 || IsMaskZ))
4546 B = Builder.CreateFNeg(
B);
4548 C = Builder.CreateFNeg(
C);
4550 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4551 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4552 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4559 if (Name.back() ==
'd')
4560 IID = Intrinsic::x86_avx512_vfmadd_f64;
4562 IID = Intrinsic::x86_avx512_vfmadd_f32;
4563 Rep = Builder.CreateIntrinsic(IID,
Ops);
4565 Rep = Builder.CreateFMA(
A,
B,
C);
4574 if (NegAcc && IsMask3)
4579 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4581 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4582 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4583 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4584 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4585 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4586 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4587 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4588 bool IsMask3 = Name[11] ==
'3';
4589 bool IsMaskZ = Name[11] ==
'z';
4591 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4592 bool NegMul = Name[2] ==
'n';
4593 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4599 if (NegMul && (IsMask3 || IsMaskZ))
4600 A = Builder.CreateFNeg(
A);
4601 if (NegMul && !(IsMask3 || IsMaskZ))
4602 B = Builder.CreateFNeg(
B);
4604 C = Builder.CreateFNeg(
C);
4611 if (Name[Name.size() - 5] ==
's')
4612 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4614 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4618 Rep = Builder.CreateFMA(
A,
B,
C);
4626 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4630 if (VecWidth == 128 && EltWidth == 32)
4631 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4632 else if (VecWidth == 256 && EltWidth == 32)
4633 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4634 else if (VecWidth == 128 && EltWidth == 64)
4635 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4636 else if (VecWidth == 256 && EltWidth == 64)
4637 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4643 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4644 Rep = Builder.CreateIntrinsic(IID,
Ops);
4645 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4646 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4647 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4648 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4649 bool IsMask3 = Name[11] ==
'3';
4650 bool IsMaskZ = Name[11] ==
'z';
4652 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4653 bool IsSubAdd = Name[3] ==
's';
4657 if (Name[Name.size() - 5] ==
's')
4658 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4660 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4665 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4667 Rep = Builder.CreateIntrinsic(IID,
Ops);
4676 Value *Odd = Builder.CreateCall(FMA,
Ops);
4677 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4678 Value *Even = Builder.CreateCall(FMA,
Ops);
4684 for (
int i = 0; i != NumElts; ++i)
4685 Idxs[i] = i + (i % 2) * NumElts;
4687 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4695 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4696 Name.starts_with(
"avx512.maskz.pternlog.")) {
4697 bool ZeroMask = Name[11] ==
'z';
4701 if (VecWidth == 128 && EltWidth == 32)
4702 IID = Intrinsic::x86_avx512_pternlog_d_128;
4703 else if (VecWidth == 256 && EltWidth == 32)
4704 IID = Intrinsic::x86_avx512_pternlog_d_256;
4705 else if (VecWidth == 512 && EltWidth == 32)
4706 IID = Intrinsic::x86_avx512_pternlog_d_512;
4707 else if (VecWidth == 128 && EltWidth == 64)
4708 IID = Intrinsic::x86_avx512_pternlog_q_128;
4709 else if (VecWidth == 256 && EltWidth == 64)
4710 IID = Intrinsic::x86_avx512_pternlog_q_256;
4711 else if (VecWidth == 512 && EltWidth == 64)
4712 IID = Intrinsic::x86_avx512_pternlog_q_512;
4718 Rep = Builder.CreateIntrinsic(IID, Args);
4722 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4723 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4724 bool ZeroMask = Name[11] ==
'z';
4725 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4728 if (VecWidth == 128 && !
High)
4729 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4730 else if (VecWidth == 256 && !
High)
4731 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4732 else if (VecWidth == 512 && !
High)
4733 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4734 else if (VecWidth == 128 &&
High)
4735 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4736 else if (VecWidth == 256 &&
High)
4737 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4738 else if (VecWidth == 512 &&
High)
4739 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4745 Rep = Builder.CreateIntrinsic(IID, Args);
4749 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4750 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4751 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4752 bool ZeroMask = Name[11] ==
'z';
4753 bool IndexForm = Name[17] ==
'i';
4755 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4756 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4757 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4758 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4759 bool ZeroMask = Name[11] ==
'z';
4760 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4763 if (VecWidth == 128 && !IsSaturating)
4764 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4765 else if (VecWidth == 256 && !IsSaturating)
4766 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4767 else if (VecWidth == 512 && !IsSaturating)
4768 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4769 else if (VecWidth == 128 && IsSaturating)
4770 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4771 else if (VecWidth == 256 && IsSaturating)
4772 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4773 else if (VecWidth == 512 && IsSaturating)
4774 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4784 if (Args[1]->
getType()->isVectorTy() &&
4787 ->isIntegerTy(32) &&
4788 Args[2]->
getType()->isVectorTy() &&
4791 ->isIntegerTy(32)) {
4792 Type *NewArgType =
nullptr;
4793 if (VecWidth == 128)
4795 else if (VecWidth == 256)
4797 else if (VecWidth == 512)
4803 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4804 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4807 Rep = Builder.CreateIntrinsic(IID, Args);
4811 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4812 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4813 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4814 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4815 bool ZeroMask = Name[11] ==
'z';
4816 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4819 if (VecWidth == 128 && !IsSaturating)
4820 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4821 else if (VecWidth == 256 && !IsSaturating)
4822 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4823 else if (VecWidth == 512 && !IsSaturating)
4824 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4825 else if (VecWidth == 128 && IsSaturating)
4826 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4827 else if (VecWidth == 256 && IsSaturating)
4828 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4829 else if (VecWidth == 512 && IsSaturating)
4830 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4840 if (Args[1]->
getType()->isVectorTy() &&
4843 ->isIntegerTy(32) &&
4844 Args[2]->
getType()->isVectorTy() &&
4847 ->isIntegerTy(32)) {
4848 Type *NewArgType =
nullptr;
4849 if (VecWidth == 128)
4851 else if (VecWidth == 256)
4853 else if (VecWidth == 512)
4859 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4860 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4863 Rep = Builder.CreateIntrinsic(IID, Args);
4867 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4868 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4869 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4871 if (Name[0] ==
'a' && Name.back() ==
'2')
4872 IID = Intrinsic::x86_addcarry_32;
4873 else if (Name[0] ==
'a' && Name.back() ==
'4')
4874 IID = Intrinsic::x86_addcarry_64;
4875 else if (Name[0] ==
's' && Name.back() ==
'2')
4876 IID = Intrinsic::x86_subborrow_32;
4877 else if (Name[0] ==
's' && Name.back() ==
'4')
4878 IID = Intrinsic::x86_subborrow_64;
4885 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4888 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4891 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4895 }
else if (Name.starts_with(
"avx512.mask.") &&
4898 }
else if (Name.starts_with(
"bmi.pdep.")) {
4900 }
else if (Name.starts_with(
"bmi.pext.")) {
4910 if (Name.starts_with(
"neon.bfcvt")) {
4911 if (Name.starts_with(
"neon.bfcvtn2")) {
4913 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4915 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4916 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4919 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4920 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4922 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4926 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4927 return Builder.CreateShuffleVector(
4930 return Builder.CreateFPTrunc(CI->
getOperand(0),
4933 }
else if (Name.starts_with(
"sve.fcvt")) {
4936 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4937 .
Case(
"sve.fcvtnt.bf16f32",
4938 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4950 if (Args[1]->
getType() != BadPredTy)
4953 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
4954 BadPredTy, Args[1]);
4955 Args[1] = Builder.CreateIntrinsic(
4956 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
4958 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
4962 if (Name ==
"neon.vcvtfp2hf")
4963 return Builder.CreateBitCast(
4964 Builder.CreateFPTrunc(
4968 if (Name ==
"neon.vcvthf2fp")
4969 return Builder.CreateFPExt(
4970 Builder.CreateBitCast(
4980 if (Name ==
"mve.vctp64.old") {
4983 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
4986 Value *C1 = Builder.CreateIntrinsic(
4987 Intrinsic::arm_mve_pred_v2i,
4989 return Builder.CreateIntrinsic(
4990 Intrinsic::arm_mve_pred_i2v,
4992 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
4993 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
4994 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
4995 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
4997 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
4998 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
4999 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5000 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5002 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5003 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5004 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5005 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5006 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5007 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5008 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5009 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5010 std::vector<Type *> Tys;
5014 case Intrinsic::arm_mve_mull_int_predicated:
5015 case Intrinsic::arm_mve_vqdmull_predicated:
5016 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5019 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5020 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5021 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5025 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5029 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5033 case Intrinsic::arm_cde_vcx1q_predicated:
5034 case Intrinsic::arm_cde_vcx1qa_predicated:
5035 case Intrinsic::arm_cde_vcx2q_predicated:
5036 case Intrinsic::arm_cde_vcx2qa_predicated:
5037 case Intrinsic::arm_cde_vcx3q_predicated:
5038 case Intrinsic::arm_cde_vcx3qa_predicated:
5045 std::vector<Value *>
Ops;
5047 Type *Ty =
Op->getType();
5048 if (Ty->getScalarSizeInBits() == 1) {
5049 Value *C1 = Builder.CreateIntrinsic(
5050 Intrinsic::arm_mve_pred_v2i,
5052 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5057 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5072 auto UpgradeLegacyWMMAIUIntrinsicCall =
5077 Args.push_back(Builder.getFalse());
5081 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5088 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5093 NewCall->copyMetadata(*CI);
5097 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5098 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5099 "intrinsic should have 7 arguments");
5102 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5104 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5105 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5106 "intrinsic should have 8 arguments");
5111 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5114 switch (
F->getIntrinsicID()) {
5117 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5118 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5119 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5120 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5121 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5122 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5137 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5140 F->getParent(),
F->getIntrinsicID(), Overloads);
5145 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5150 NewCall->copyMetadata(*CI);
5151 NewCall->takeName(CI);
5173 if (NumOperands < 3)
5186 bool IsVolatile =
false;
5190 if (NumOperands > 3)
5195 if (NumOperands > 5) {
5197 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5211 if (VT->getElementType()->isIntegerTy(16)) {
5214 Val = Builder.CreateBitCast(Val, AsBF16);
5222 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5224 unsigned AddrSpace = PtrTy->getAddressSpace();
5227 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5229 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5234 MDNode *RangeNotPrivate =
5237 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5243 return Builder.CreateBitCast(RMW, RetTy);
5264 return MAV->getMetadata();
5273 if (Name ==
"label") {
5275 }
else if (Name ==
"assign") {
5282 }
else if (Name ==
"declare") {
5286 }
else if (Name ==
"addr") {
5296 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5297 }
else if (Name ==
"value") {
5300 unsigned ExprOp = 2;
5315 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5323 int64_t OffsetVal =
Offset->getSExtValue();
5324 return Builder.CreateIntrinsic(OffsetVal >= 0
5325 ? Intrinsic::vector_splice_left
5326 : Intrinsic::vector_splice_right,
5328 {CI->getArgOperand(0), CI->getArgOperand(1),
5329 Builder.getInt32(std::abs(OffsetVal))});
5334 if (Name.starts_with(
"to.fp16")) {
5336 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5337 return Builder.CreateBitCast(Cast, CI->
getType());
5340 if (Name.starts_with(
"from.fp16")) {
5342 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5343 return Builder.CreateFPExt(Cast, CI->
getType());
5402 else if (Opcode == Instruction::ICmp)
5405 else if (Opcode == Instruction::FCmp)
5408 else if (Opcode == Instruction::Select)
5413 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5425 if (Defaults.empty())
5428 unsigned OldArgCount = CI->
arg_size();
5429 unsigned NewArgCount = NewFn->
arg_size();
5433 if (OldArgCount >= NewArgCount)
5441 if (OldArgCount < FirstDefault)
5446 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5447 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5448 "missing argument outside the default range");
5449 Type *ParamTy = NewFT->getParamType(Idx);
5454 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5460 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5492 if (!Name.consume_front(
"llvm."))
5495 bool IsX86 = Name.consume_front(
"x86.");
5496 bool IsNVVM = Name.consume_front(
"nvvm.");
5497 bool IsAArch64 = Name.consume_front(
"aarch64.");
5498 bool IsARM = Name.consume_front(
"arm.");
5499 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5500 bool IsDbg = Name.consume_front(
"dbg.");
5502 (Name.consume_front(
"experimental.vector.splice") ||
5503 Name.consume_front(
"vector.splice")) &&
5504 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5505 Value *Rep =
nullptr;
5507 if (!IsX86 && Name ==
"stackprotectorcheck") {
5509 }
else if (IsNVVM) {
5513 }
else if (IsAArch64) {
5517 }
else if (IsAMDGCN) {
5521 }
else if (IsOldSplice) {
5523 }
else if (Name.consume_front(
"convert.")) {
5525 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5540 const auto &DefaultCase = [&]() ->
void {
5548 "Unknown function for CallBase upgrade and isn't just a name change");
5556 "Return type must have changed");
5557 assert(OldST->getNumElements() ==
5559 "Must have same number of elements");
5562 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5565 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5566 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5567 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5591 case Intrinsic::arm_neon_vst1:
5592 case Intrinsic::arm_neon_vst2:
5593 case Intrinsic::arm_neon_vst3:
5594 case Intrinsic::arm_neon_vst4:
5595 case Intrinsic::arm_neon_vst2lane:
5596 case Intrinsic::arm_neon_vst3lane:
5597 case Intrinsic::arm_neon_vst4lane: {
5599 NewCall = Builder.CreateCall(NewFn, Args);
5602 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5603 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5604 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5609 NewCall = Builder.CreateCall(NewFn, Args);
5612 case Intrinsic::aarch64_sve_ld3_sret:
5613 case Intrinsic::aarch64_sve_ld4_sret:
5614 case Intrinsic::aarch64_sve_ld2_sret: {
5622 Name = Name.substr(5);
5629 unsigned MinElts = RetTy->getMinNumElements() /
N;
5631 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5633 for (
unsigned I = 0;
I <
N;
I++) {
5634 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5635 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5641 case Intrinsic::coro_end_async:
5642 case Intrinsic::coro_end: {
5644 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5646 NewCall = Builder.CreateCall(NewFn, Args);
5651 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5652 Value *InRamp = Builder.CreateCall(IsInRamp);
5662 case Intrinsic::vector_extract: {
5664 Name = Name.substr(5);
5665 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5670 unsigned MinElts = RetTy->getMinNumElements();
5673 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5677 case Intrinsic::vector_insert: {
5679 Name = Name.substr(5);
5680 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5684 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5689 NewCall = Builder.CreateCall(
5693 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5699 assert(
N > 1 &&
"Create is expected to be between 2-4");
5702 unsigned MinElts = RetTy->getMinNumElements() /
N;
5703 for (
unsigned I = 0;
I <
N;
I++) {
5705 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5712 case Intrinsic::arm_neon_bfdot:
5713 case Intrinsic::arm_neon_bfmmla:
5714 case Intrinsic::arm_neon_bfmlalb:
5715 case Intrinsic::arm_neon_bfmlalt:
5716 case Intrinsic::aarch64_neon_bfdot:
5717 case Intrinsic::aarch64_neon_bfmmla:
5718 case Intrinsic::aarch64_neon_bfmlalb:
5719 case Intrinsic::aarch64_neon_bfmlalt: {
5722 "Mismatch between function args and call args");
5723 size_t OperandWidth =
5725 assert((OperandWidth == 64 || OperandWidth == 128) &&
5726 "Unexpected operand width");
5728 auto Iter = CI->
args().begin();
5729 Args.push_back(*Iter++);
5730 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5731 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5732 NewCall = Builder.CreateCall(NewFn, Args);
5736 case Intrinsic::bitreverse:
5737 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5740 case Intrinsic::ctlz:
5741 case Intrinsic::cttz: {
5748 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5752 case Intrinsic::objectsize: {
5753 Value *NullIsUnknownSize =
5757 NewCall = Builder.CreateCall(
5762 case Intrinsic::ctpop:
5763 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5765 case Intrinsic::dbg_value: {
5767 Name = Name.substr(5);
5769 if (Name.starts_with(
"dbg.addr")) {
5783 if (
Offset->isNullValue()) {
5784 NewCall = Builder.CreateCall(
5793 case Intrinsic::ptr_annotation:
5801 NewCall = Builder.CreateCall(
5810 case Intrinsic::var_annotation:
5817 NewCall = Builder.CreateCall(
5826 case Intrinsic::riscv_aes32dsi:
5827 case Intrinsic::riscv_aes32dsmi:
5828 case Intrinsic::riscv_aes32esi:
5829 case Intrinsic::riscv_aes32esmi:
5830 case Intrinsic::riscv_sm4ks:
5831 case Intrinsic::riscv_sm4ed: {
5841 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5842 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5848 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5849 Value *Res = NewCall;
5851 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5857 case Intrinsic::nvvm_mapa_shared_cluster: {
5861 Value *Res = NewCall;
5862 Res = Builder.CreateAddrSpaceCast(
5869 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5870 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5873 Args[0] = Builder.CreateAddrSpaceCast(
5876 NewCall = Builder.CreateCall(NewFn, Args);
5882 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5883 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5884 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5885 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5886 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5887 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5888 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5889 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5896 Args[0] = Builder.CreateAddrSpaceCast(
5905 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5907 NewCall = Builder.CreateCall(NewFn, Args);
5913 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
5914 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
5915 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
5916 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
5917 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
5918 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
5919 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
5920 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
5922 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
5926 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
5927 NewCall = Builder.CreateCall(NewFn, Args);
5930 case Intrinsic::nvvm_tcgen05_mma_shared:
5931 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5932 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5933 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
5934 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
5935 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
5936 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
5937 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
5938 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
5939 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
5940 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
5941 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
5942 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
5943 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
5944 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
5945 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
5946 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
5947 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
5948 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
5949 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
5950 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
5951 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
5952 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
5953 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
5954 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
5955 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
5956 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
5957 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
5958 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
5959 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
5960 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
5961 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
5962 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
5963 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
5964 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
5965 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
5966 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
5967 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
5968 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
5970 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
5971 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
5973 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
5974 case Intrinsic::nvvm_tcgen05_mma_tensor:
5975 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
5976 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
5977 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
5978 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
5979 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
5980 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
5981 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
5982 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
5983 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
5984 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
5985 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
5986 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
5987 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
5988 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
5990 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
5991 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
5993 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
5995 Args.push_back(Builder.getInt32(0));
5996 NewCall = Builder.CreateCall(NewFn, Args);
5999 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6000 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6001 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6002 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6005 Builder.getFalse()});
6007 case Intrinsic::riscv_sha256sig0:
6008 case Intrinsic::riscv_sha256sig1:
6009 case Intrinsic::riscv_sha256sum0:
6010 case Intrinsic::riscv_sha256sum1:
6011 case Intrinsic::riscv_sm3p0:
6012 case Intrinsic::riscv_sm3p1: {
6019 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6021 NewCall = Builder.CreateCall(NewFn, Arg);
6023 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6030 case Intrinsic::x86_xop_vfrcz_ss:
6031 case Intrinsic::x86_xop_vfrcz_sd:
6032 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6035 case Intrinsic::x86_xop_vpermil2pd:
6036 case Intrinsic::x86_xop_vpermil2ps:
6037 case Intrinsic::x86_xop_vpermil2pd_256:
6038 case Intrinsic::x86_xop_vpermil2ps_256: {
6042 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6043 NewCall = Builder.CreateCall(NewFn, Args);
6047 case Intrinsic::x86_sse41_ptestc:
6048 case Intrinsic::x86_sse41_ptestz:
6049 case Intrinsic::x86_sse41_ptestnzc: {
6063 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6064 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6066 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6070 case Intrinsic::x86_rdtscp: {
6076 NewCall = Builder.CreateCall(NewFn);
6078 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6081 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6089 case Intrinsic::x86_sse41_insertps:
6090 case Intrinsic::x86_sse41_dppd:
6091 case Intrinsic::x86_sse41_dpps:
6092 case Intrinsic::x86_sse41_mpsadbw:
6093 case Intrinsic::x86_avx_dp_ps_256:
6094 case Intrinsic::x86_avx2_mpsadbw: {
6100 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6101 NewCall = Builder.CreateCall(NewFn, Args);
6105 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6106 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6107 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6108 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6109 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6110 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6116 NewCall = Builder.CreateCall(NewFn, Args);
6125 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6126 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6127 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6128 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6129 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6130 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6134 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6135 Args[1] = Builder.CreateBitCast(
6138 NewCall = Builder.CreateCall(NewFn, Args);
6139 Value *Res = Builder.CreateBitCast(
6147 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6148 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6149 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6153 Args[1] = Builder.CreateBitCast(
6155 Args[2] = Builder.CreateBitCast(
6158 NewCall = Builder.CreateCall(NewFn, Args);
6162 case Intrinsic::thread_pointer: {
6163 NewCall = Builder.CreateCall(NewFn, {});
6167 case Intrinsic::memcpy:
6168 case Intrinsic::memmove:
6169 case Intrinsic::memset: {
6185 NewCall = Builder.CreateCall(NewFn, Args);
6187 AttributeList NewAttrs = AttributeList::get(
6188 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6189 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6190 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6195 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6198 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6202 case Intrinsic::masked_load:
6203 case Intrinsic::masked_gather:
6204 case Intrinsic::masked_store:
6205 case Intrinsic::masked_scatter: {
6211 auto GetMaybeAlign = [](
Value *
Op) {
6213 uint64_t Val = CI->getZExtValue();
6221 auto GetAlign = [&](
Value *
Op) {
6230 case Intrinsic::masked_load:
6231 NewCall = Builder.CreateMaskedLoad(
6235 case Intrinsic::masked_gather:
6236 NewCall = Builder.CreateMaskedGather(
6242 case Intrinsic::masked_store:
6243 NewCall = Builder.CreateMaskedStore(
6247 case Intrinsic::masked_scatter:
6248 NewCall = Builder.CreateMaskedScatter(
6250 DL.getValueOrABITypeAlignment(
6264 case Intrinsic::lifetime_start:
6265 case Intrinsic::lifetime_end: {
6277 NewCall = Builder.CreateLifetimeStart(Ptr);
6279 NewCall = Builder.CreateLifetimeEnd(Ptr);
6288 case Intrinsic::x86_avx512_vpdpbusd_128:
6289 case Intrinsic::x86_avx512_vpdpbusd_256:
6290 case Intrinsic::x86_avx512_vpdpbusd_512:
6291 case Intrinsic::x86_avx512_vpdpbusds_128:
6292 case Intrinsic::x86_avx512_vpdpbusds_256:
6293 case Intrinsic::x86_avx512_vpdpbusds_512:
6294 case Intrinsic::x86_avx2_vpdpbssd_128:
6295 case Intrinsic::x86_avx2_vpdpbssd_256:
6296 case Intrinsic::x86_avx10_vpdpbssd_512:
6297 case Intrinsic::x86_avx2_vpdpbssds_128:
6298 case Intrinsic::x86_avx2_vpdpbssds_256:
6299 case Intrinsic::x86_avx10_vpdpbssds_512:
6300 case Intrinsic::x86_avx2_vpdpbsud_128:
6301 case Intrinsic::x86_avx2_vpdpbsud_256:
6302 case Intrinsic::x86_avx10_vpdpbsud_512:
6303 case Intrinsic::x86_avx2_vpdpbsuds_128:
6304 case Intrinsic::x86_avx2_vpdpbsuds_256:
6305 case Intrinsic::x86_avx10_vpdpbsuds_512:
6306 case Intrinsic::x86_avx2_vpdpbuud_128:
6307 case Intrinsic::x86_avx2_vpdpbuud_256:
6308 case Intrinsic::x86_avx10_vpdpbuud_512:
6309 case Intrinsic::x86_avx2_vpdpbuuds_128:
6310 case Intrinsic::x86_avx2_vpdpbuuds_256:
6311 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6316 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6317 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6319 NewCall = Builder.CreateCall(NewFn, Args);
6322 case Intrinsic::x86_avx512_vpdpwssd_128:
6323 case Intrinsic::x86_avx512_vpdpwssd_256:
6324 case Intrinsic::x86_avx512_vpdpwssd_512:
6325 case Intrinsic::x86_avx512_vpdpwssds_128:
6326 case Intrinsic::x86_avx512_vpdpwssds_256:
6327 case Intrinsic::x86_avx512_vpdpwssds_512:
6328 case Intrinsic::x86_avx2_vpdpwsud_128:
6329 case Intrinsic::x86_avx2_vpdpwsud_256:
6330 case Intrinsic::x86_avx10_vpdpwsud_512:
6331 case Intrinsic::x86_avx2_vpdpwsuds_128:
6332 case Intrinsic::x86_avx2_vpdpwsuds_256:
6333 case Intrinsic::x86_avx10_vpdpwsuds_512:
6334 case Intrinsic::x86_avx2_vpdpwusd_128:
6335 case Intrinsic::x86_avx2_vpdpwusd_256:
6336 case Intrinsic::x86_avx10_vpdpwusd_512:
6337 case Intrinsic::x86_avx2_vpdpwusds_128:
6338 case Intrinsic::x86_avx2_vpdpwusds_256:
6339 case Intrinsic::x86_avx10_vpdpwusds_512:
6340 case Intrinsic::x86_avx2_vpdpwuud_128:
6341 case Intrinsic::x86_avx2_vpdpwuud_256:
6342 case Intrinsic::x86_avx10_vpdpwuud_512:
6343 case Intrinsic::x86_avx2_vpdpwuuds_128:
6344 case Intrinsic::x86_avx2_vpdpwuuds_256:
6345 case Intrinsic::x86_avx10_vpdpwuuds_512:
6350 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6351 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6353 NewCall = Builder.CreateCall(NewFn, Args);
6356 assert(NewCall &&
"Should have either set this variable or returned through "
6357 "the default case");
6364 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6378 F->eraseFromParent();
6384 if (NumOperands == 0)
6392 if (NumOperands == 3) {
6396 Metadata *Elts2[] = {ScalarType, ScalarType,
6410 if (
Opc != Instruction::BitCast)
6414 Type *SrcTy = V->getType();
6431 if (
Opc != Instruction::BitCast)
6434 Type *SrcTy =
C->getType();
6451 if (Flag.getNumOperands() < 3)
6452 return std::nullopt;
6454 return Name->getString();
6455 return std::nullopt;
6469 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6470 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6471 if (auto Name = getModuleFlagNameSafely(*Flag))
6472 return *Name ==
"Debug Info Version";
6475 if (OpIt != ModFlags->op_end()) {
6476 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6483 bool BrokenDebugInfo =
false;
6486 if (!BrokenDebugInfo)
6492 M.getContext().diagnose(Diag);
6499 M.getContext().diagnose(DiagVersion);
6509 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6512 if (
F->hasFnAttribute(Attr)) {
6515 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6517 auto [Part, Rest] = S.
split(
',');
6523 const unsigned Dim = DimC -
'x';
6524 assert(Dim < 3 &&
"Unexpected dim char");
6534 F->addFnAttr(Attr, NewAttr);
6538 return S ==
"x" || S ==
"y" || S ==
"z";
6543 if (K ==
"kernel") {
6555 const unsigned Idx = (AlignIdxValuePair >> 16);
6556 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6561 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6566 if (K ==
"minctasm") {
6571 if (K ==
"maxnreg") {
6576 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6580 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6584 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6588 if (K ==
"grid_constant") {
6603 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6610 if (!SeenNodes.
insert(MD).second)
6617 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6624 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6626 const MDOperand &V = MD->getOperand(j + 1);
6629 NewOperands.
append({K, V});
6632 if (NewOperands.
size() > 1)
6645 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6646 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6647 if (ModRetainReleaseMarker) {
6653 ID->getString().split(ValueComp,
"#");
6654 if (ValueComp.
size() == 2) {
6655 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6659 M.eraseNamedMetadata(ModRetainReleaseMarker);
6670 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6696 bool InvalidCast =
false;
6698 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6711 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6713 Args.push_back(Arg);
6720 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6725 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6738 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6746 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6747 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6748 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6749 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6750 {
"objc_autoreleaseReturnValue",
6751 llvm::Intrinsic::objc_autoreleaseReturnValue},
6752 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6753 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6754 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6755 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6756 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6757 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6758 {
"objc_release", llvm::Intrinsic::objc_release},
6759 {
"objc_retain", llvm::Intrinsic::objc_retain},
6760 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6761 {
"objc_retainAutoreleaseReturnValue",
6762 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6763 {
"objc_retainAutoreleasedReturnValue",
6764 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6765 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6766 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6767 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6768 {
"objc_unsafeClaimAutoreleasedReturnValue",
6769 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6770 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6771 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6772 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6773 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6774 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6775 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6776 {
"objc_arc_annotation_topdown_bbstart",
6777 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6778 {
"objc_arc_annotation_topdown_bbend",
6779 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6780 {
"objc_arc_annotation_bottomup_bbstart",
6781 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6782 {
"objc_arc_annotation_bottomup_bbend",
6783 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6785 for (
auto &
I : RuntimeFuncs)
6786 UpgradeToIntrinsic(
I.first,
I.second);
6810 std::optional<bool> UseAddressDisc;
6813 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6814 for (
const MDNode *Flag : ModFlags->operands()) {
6816 if (Name && (*Name ==
"ptrauth-init-fini" ||
6817 *Name ==
"ptrauth-init-fini-address-discrimination"))
6822 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6823 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6824 constexpr unsigned ExpectedAddressMarker = 1;
6827 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6830 bool HasAddressDisc;
6831 if (!CPA->hasAddressDiscriminator())
6832 HasAddressDisc =
false;
6833 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6834 HasAddressDisc =
true;
6838 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6841 UseAddressDisc = HasAddressDisc;
6842 return CPA->getPointer();
6846 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6849 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6851 if (!GV || !GV->hasInitializer())
6855 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6858 std::vector<Constant *> NewStructors;
6859 NewStructors.reserve(OldStructorsArray->getNumOperands());
6861 for (
Use &U : OldStructorsArray->operands()) {
6870 Func = UpgradeSinglePointer(Func);
6874 NewStructors.push_back(
6883 if (GlobalArraysToUpgrade.
empty())
6885 assert(UseAddressDisc.has_value());
6887 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6888 GV->setInitializer(NewInit);
6891 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6901 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6905 bool HasObjCFlag =
false, HasClassProperties =
false;
6906 bool HasSwiftVersionFlag =
false;
6907 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6914 if (
Op->getNumOperands() != 3)
6928 if (ID->getString() ==
"Objective-C Image Info Version")
6930 if (ID->getString() ==
"Objective-C Class Properties")
6931 HasClassProperties =
true;
6933 if (ID->getString() ==
"PIC Level") {
6934 if (
auto *Behavior =
6936 uint64_t V = Behavior->getLimitedValue();
6942 if (ID->getString() ==
"PIE Level")
6943 if (
auto *Behavior =
6950 if (ID->getString() ==
"branch-target-enforcement" ||
6951 ID->getString().starts_with(
"sign-return-address")) {
6952 if (
auto *Behavior =
6958 Op->getOperand(1),
Op->getOperand(2)};
6968 if (ID->getString() ==
"Objective-C Image Info Section") {
6971 Value->getString().split(ValueComp,
" ");
6972 if (ValueComp.
size() != 1) {
6973 std::string NewValue;
6974 for (
auto &S : ValueComp)
6975 NewValue += S.str();
6986 if (ID->getString() ==
"Objective-C Garbage Collection") {
6989 assert(Md->getValue() &&
"Expected non-empty metadata");
6990 auto Type = Md->getValue()->getType();
6993 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
6994 if ((Val & 0xff) != Val) {
6995 HasSwiftVersionFlag =
true;
6996 SwiftABIVersion = (Val & 0xff00) >> 8;
6997 SwiftMajorVersion = (Val & 0xff000000) >> 24;
6998 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7009 if (ID->getString() ==
"amdgpu_code_object_version") {
7012 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7021 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7050 if (HasObjCFlag && !HasClassProperties) {
7056 if (HasSwiftVersionFlag) {
7060 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7062 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7070 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7074 auto MatchesVersion = [](
const MDNode *
Op) {
7075 return Op->getNumOperands() >= 3 &&
7089 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7090 assert(
Op->getNumOperands() >= 2 &&
7091 "Expected at least 2 operands - name and linkage type");
7103 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7114 auto TrimSpaces = [](
StringRef Section) -> std::string {
7116 Section.split(Components,
',');
7121 for (
auto Component : Components)
7122 OS <<
',' << Component.trim();
7127 for (
auto &GV : M.globals()) {
7128 if (!GV.hasSection())
7133 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7138 GV.setSection(TrimSpaces(Section));
7154struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7155 StrictFPUpgradeVisitor() =
default;
7158 if (!
Call.isStrictFP())
7164 Call.removeFnAttr(Attribute::StrictFP);
7165 Call.addFnAttr(Attribute::NoBuiltin);
7170struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7171 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7172 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7174 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7189 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7190 StrictFPUpgradeVisitor SFPV;
7195 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7196 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7197 for (
auto &Arg :
F.args())
7199 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7201 bool AddingAttrs =
false, RemovingAttrs =
false;
7202 AttrBuilder AttrsToAdd(
F.getContext());
7207 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7208 A.isValid() &&
A.isStringAttribute()) {
7209 F.setSection(
A.getValueAsString());
7211 RemovingAttrs =
true;
7215 A.isValid() &&
A.isStringAttribute()) {
7218 AddingAttrs = RemovingAttrs =
true;
7221 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7222 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7224 RemovingAttrs =
true;
7225 if (
A.getValueAsString() ==
"true") {
7226 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7235 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7238 if (
A.getValueAsBool()) {
7239 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7245 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7246 RemovingAttrs =
true;
7253 bool HandleDenormalMode =
false;
7255 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7258 DenormalFPMath = ParsedMode;
7260 AddingAttrs = RemovingAttrs =
true;
7261 HandleDenormalMode =
true;
7265 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7269 DenormalFPMathF32 = ParsedMode;
7271 AddingAttrs = RemovingAttrs =
true;
7272 HandleDenormalMode =
true;
7276 if (HandleDenormalMode)
7277 AttrsToAdd.addDenormalFPEnvAttr(
7281 F.removeFnAttrs(AttrsToRemove);
7284 F.addFnAttrs(AttrsToAdd);
7290 if (!
F.hasFnAttribute(FnAttrName))
7291 F.addFnAttr(FnAttrName,
Value);
7298 if (!
F.hasFnAttribute(FnAttrName)) {
7300 F.addFnAttr(FnAttrName);
7302 auto A =
F.getFnAttribute(FnAttrName);
7303 if (
"false" ==
A.getValueAsString())
7304 F.removeFnAttr(FnAttrName);
7305 else if (
"true" ==
A.getValueAsString()) {
7306 F.removeFnAttr(FnAttrName);
7307 F.addFnAttr(FnAttrName);
7313 Triple T(M.getTargetTriple());
7314 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7317 uint64_t BTEValue = 0;
7318 uint64_t BPPLRValue = 0;
7319 uint64_t GCSValue = 0;
7320 uint64_t SRAValue = 0;
7321 uint64_t SRAALLValue = 0;
7322 uint64_t SRABKeyValue = 0;
7324 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7328 if (
Op->getNumOperands() != 3)
7337 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7338 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7339 : IDStr ==
"guarded-control-stack" ? &GCSValue
7340 : IDStr ==
"sign-return-address" ? &SRAValue
7341 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7342 : IDStr ==
"sign-return-address-with-bkey"
7348 *ValPtr = CI->getZExtValue();
7354 bool BTE = BTEValue == 1;
7355 bool BPPLR = BPPLRValue == 1;
7356 bool GCS = GCSValue == 1;
7357 bool SRA = SRAValue == 1;
7360 if (SRA && SRAALLValue == 1)
7361 SignTypeValue =
"all";
7364 if (SRA && SRABKeyValue == 1)
7365 SignKeyValue =
"b_key";
7367 for (
Function &
F : M.getFunctionList()) {
7368 if (
F.isDeclaration())
7375 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7376 A.isValid() &&
"none" ==
A.getValueAsString()) {
7377 F.removeFnAttr(
"sign-return-address");
7378 F.removeFnAttr(
"sign-return-address-key");
7394 if (SRAALLValue == 1)
7396 if (SRABKeyValue == 1)
7423 if (
T->getNumOperands() < 1)
7428 if (S->getString().starts_with(
"llvm.vectorizer."))
7434 StringRef OldPrefix =
"llvm.vectorizer.";
7437 if (OldTag ==
"llvm.vectorizer.unroll")
7449 if (
T->getNumOperands() < 1)
7461 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7474 Ops.reserve(
T->getNumOperands());
7475 Ops.push_back(NewTag);
7476 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7477 Ops.push_back(
T->getOperand(
I));
7494 if (
T->isDistinct()) {
7495 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7507 Ops.reserve(
T->getNumOperands());
7518 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7519 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7520 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7523 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7525 auto I =
DL.find(
"-n64-");
7527 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7532 std::string Res =
DL.str();
7535 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7536 Res.append(Res.empty() ?
"G1" :
"-G1");
7544 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7545 Res.append(
"-ni:7:8:9");
7547 if (
DL.ends_with(
"ni:7"))
7549 if (
DL.ends_with(
"ni:7:8"))
7554 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7555 Res.append(
"-p7:160:256:256:32");
7556 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7557 Res.append(
"-p8:128:128:128:48");
7558 constexpr StringRef OldP8(
"-p8:128:128-");
7559 if (
DL.contains(OldP8))
7560 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7561 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7562 Res.append(
"-p9:192:256:256:32");
7566 if (!
DL.contains(
"m:e"))
7567 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7572 if (
T.isSystemZ() && !
DL.empty()) {
7574 if (!
DL.contains(
"-S64"))
7575 return "E-S64" +
DL.drop_front(1).str();
7579 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7582 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7583 if (!
DL.contains(AddrSpaces)) {
7585 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7586 if (R.match(Res, &
Groups))
7592 if (
T.isAArch64()) {
7594 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7595 Res.append(
"-Fn32");
7596 AddPtr32Ptr64AddrSpaces();
7600 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7604 std::string I64 =
"-i64:64";
7605 std::string I128 =
"-i128:128";
7607 size_t Pos = Res.find(I64);
7608 if (Pos !=
size_t(-1))
7609 Res.insert(Pos + I64.size(), I128);
7613 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7614 size_t Pos = Res.find(
"-S128");
7617 Res.insert(Pos,
"-f64:32:64");
7623 AddPtr32Ptr64AddrSpaces();
7631 if (!
T.isOSIAMCU()) {
7632 std::string I128 =
"-i128:128";
7635 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7636 if (R.match(Res, &
Groups))
7644 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7646 auto I =
Ref.find(
"-f80:32-");
7648 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7656 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7659 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7660 B.removeAttribute(
"no-frame-pointer-elim");
7662 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7664 if (FramePointer !=
"all")
7665 FramePointer =
"non-leaf";
7666 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7668 if (!FramePointer.
empty())
7669 B.addAttribute(
"frame-pointer", FramePointer);
7671 A =
B.getAttribute(
"null-pointer-is-valid");
7674 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7675 B.removeAttribute(
"null-pointer-is-valid");
7676 if (NullPointerIsValid)
7677 B.addAttribute(Attribute::NullPointerIsValid);
7680 A =
B.getAttribute(
"uniform-work-group-size");
7684 bool IsTrue = Val ==
"true";
7685 B.removeAttribute(
"uniform-work-group-size");
7687 B.addAttribute(
"uniform-work-group-size");
7698 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.