35#include "llvm/IR/IntrinsicsAArch64.h"
36#include "llvm/IR/IntrinsicsAMDGPU.h"
37#include "llvm/IR/IntrinsicsARM.h"
38#include "llvm/IR/IntrinsicsNVPTX.h"
39#include "llvm/IR/IntrinsicsRISCV.h"
40#include "llvm/IR/IntrinsicsWebAssembly.h"
41#include "llvm/IR/IntrinsicsX86.h"
66 cl::desc(
"Disable autoupgrade of debug info"));
85 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
100 Type *LastArgType =
F->getFunctionType()->getParamType(
101 F->getFunctionType()->getNumParams() - 1);
116 if (
F->getReturnType()->isVectorTy())
129 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
130 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
147 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
148 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
162 if (
F->getReturnType()->getScalarType()->isBFloatTy())
172 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
186 if (Name.consume_front(
"avx."))
187 return (Name.starts_with(
"blend.p") ||
188 Name ==
"cvt.ps2.pd.256" ||
189 Name ==
"cvtdq2.pd.256" ||
190 Name ==
"cvtdq2.ps.256" ||
191 Name.starts_with(
"movnt.") ||
192 Name.starts_with(
"sqrt.p") ||
193 Name.starts_with(
"storeu.") ||
194 Name.starts_with(
"vbroadcast.s") ||
195 Name.starts_with(
"vbroadcastf128") ||
196 Name.starts_with(
"vextractf128.") ||
197 Name.starts_with(
"vinsertf128.") ||
198 Name.starts_with(
"vperm2f128.") ||
199 Name.starts_with(
"vpermil."));
201 if (Name.consume_front(
"avx2."))
202 return (Name ==
"movntdqa" ||
203 Name.starts_with(
"pabs.") ||
204 Name.starts_with(
"padds.") ||
205 Name.starts_with(
"paddus.") ||
206 Name.starts_with(
"pblendd.") ||
208 Name.starts_with(
"pbroadcast") ||
209 Name.starts_with(
"pcmpeq.") ||
210 Name.starts_with(
"pcmpgt.") ||
211 Name.starts_with(
"pmax") ||
212 Name.starts_with(
"pmin") ||
213 Name.starts_with(
"pmovsx") ||
214 Name.starts_with(
"pmovzx") ||
216 Name ==
"pmulu.dq" ||
217 Name.starts_with(
"psll.dq") ||
218 Name.starts_with(
"psrl.dq") ||
219 Name.starts_with(
"psubs.") ||
220 Name.starts_with(
"psubus.") ||
221 Name.starts_with(
"vbroadcast") ||
222 Name ==
"vbroadcasti128" ||
223 Name ==
"vextracti128" ||
224 Name ==
"vinserti128" ||
225 Name ==
"vperm2i128");
227 if (Name.consume_front(
"avx512.")) {
228 if (Name.consume_front(
"mask."))
230 return (Name.starts_with(
"add.p") ||
231 Name.starts_with(
"and.") ||
232 Name.starts_with(
"andn.") ||
233 Name.starts_with(
"broadcast.s") ||
234 Name.starts_with(
"broadcastf32x4.") ||
235 Name.starts_with(
"broadcastf32x8.") ||
236 Name.starts_with(
"broadcastf64x2.") ||
237 Name.starts_with(
"broadcastf64x4.") ||
238 Name.starts_with(
"broadcasti32x4.") ||
239 Name.starts_with(
"broadcasti32x8.") ||
240 Name.starts_with(
"broadcasti64x2.") ||
241 Name.starts_with(
"broadcasti64x4.") ||
242 Name.starts_with(
"cmp.b") ||
243 Name.starts_with(
"cmp.d") ||
244 Name.starts_with(
"cmp.q") ||
245 Name.starts_with(
"cmp.w") ||
246 Name.starts_with(
"compress.b") ||
247 Name.starts_with(
"compress.d") ||
248 Name.starts_with(
"compress.p") ||
249 Name.starts_with(
"compress.q") ||
250 Name.starts_with(
"compress.store.") ||
251 Name.starts_with(
"compress.w") ||
252 Name.starts_with(
"conflict.") ||
253 Name.starts_with(
"cvtdq2pd.") ||
254 Name.starts_with(
"cvtdq2ps.") ||
255 Name ==
"cvtpd2dq.256" ||
256 Name ==
"cvtpd2ps.256" ||
257 Name ==
"cvtps2pd.128" ||
258 Name ==
"cvtps2pd.256" ||
259 Name.starts_with(
"cvtqq2pd.") ||
260 Name ==
"cvtqq2ps.256" ||
261 Name ==
"cvtqq2ps.512" ||
262 Name ==
"cvttpd2dq.256" ||
263 Name ==
"cvttps2dq.128" ||
264 Name ==
"cvttps2dq.256" ||
265 Name.starts_with(
"cvtudq2pd.") ||
266 Name.starts_with(
"cvtudq2ps.") ||
267 Name.starts_with(
"cvtuqq2pd.") ||
268 Name ==
"cvtuqq2ps.256" ||
269 Name ==
"cvtuqq2ps.512" ||
270 Name.starts_with(
"dbpsadbw.") ||
271 Name.starts_with(
"div.p") ||
272 Name.starts_with(
"expand.b") ||
273 Name.starts_with(
"expand.d") ||
274 Name.starts_with(
"expand.load.") ||
275 Name.starts_with(
"expand.p") ||
276 Name.starts_with(
"expand.q") ||
277 Name.starts_with(
"expand.w") ||
278 Name.starts_with(
"fpclass.p") ||
279 Name.starts_with(
"insert") ||
280 Name.starts_with(
"load.") ||
281 Name.starts_with(
"loadu.") ||
282 Name.starts_with(
"lzcnt.") ||
283 Name.starts_with(
"max.p") ||
284 Name.starts_with(
"min.p") ||
285 Name.starts_with(
"movddup") ||
286 Name.starts_with(
"move.s") ||
287 Name.starts_with(
"movshdup") ||
288 Name.starts_with(
"movsldup") ||
289 Name.starts_with(
"mul.p") ||
290 Name.starts_with(
"or.") ||
291 Name.starts_with(
"pabs.") ||
292 Name.starts_with(
"packssdw.") ||
293 Name.starts_with(
"packsswb.") ||
294 Name.starts_with(
"packusdw.") ||
295 Name.starts_with(
"packuswb.") ||
296 Name.starts_with(
"padd.") ||
297 Name.starts_with(
"padds.") ||
298 Name.starts_with(
"paddus.") ||
299 Name.starts_with(
"palignr.") ||
300 Name.starts_with(
"pand.") ||
301 Name.starts_with(
"pandn.") ||
302 Name.starts_with(
"pavg") ||
303 Name.starts_with(
"pbroadcast") ||
304 Name.starts_with(
"pcmpeq.") ||
305 Name.starts_with(
"pcmpgt.") ||
306 Name.starts_with(
"perm.df.") ||
307 Name.starts_with(
"perm.di.") ||
308 Name.starts_with(
"permvar.") ||
309 Name.starts_with(
"pmaddubs.w.") ||
310 Name.starts_with(
"pmaddw.d.") ||
311 Name.starts_with(
"pmax") ||
312 Name.starts_with(
"pmin") ||
313 Name ==
"pmov.qd.256" ||
314 Name ==
"pmov.qd.512" ||
315 Name ==
"pmov.wb.256" ||
316 Name ==
"pmov.wb.512" ||
317 Name.starts_with(
"pmovsx") ||
318 Name.starts_with(
"pmovzx") ||
319 Name.starts_with(
"pmul.dq.") ||
320 Name.starts_with(
"pmul.hr.sw.") ||
321 Name.starts_with(
"pmulh.w.") ||
322 Name.starts_with(
"pmulhu.w.") ||
323 Name.starts_with(
"pmull.") ||
324 Name.starts_with(
"pmultishift.qb.") ||
325 Name.starts_with(
"pmulu.dq.") ||
326 Name.starts_with(
"por.") ||
327 Name.starts_with(
"prol.") ||
328 Name.starts_with(
"prolv.") ||
329 Name.starts_with(
"pror.") ||
330 Name.starts_with(
"prorv.") ||
331 Name.starts_with(
"pshuf.b.") ||
332 Name.starts_with(
"pshuf.d.") ||
333 Name.starts_with(
"pshufh.w.") ||
334 Name.starts_with(
"pshufl.w.") ||
335 Name.starts_with(
"psll.d") ||
336 Name.starts_with(
"psll.q") ||
337 Name.starts_with(
"psll.w") ||
338 Name.starts_with(
"pslli") ||
339 Name.starts_with(
"psllv") ||
340 Name.starts_with(
"psra.d") ||
341 Name.starts_with(
"psra.q") ||
342 Name.starts_with(
"psra.w") ||
343 Name.starts_with(
"psrai") ||
344 Name.starts_with(
"psrav") ||
345 Name.starts_with(
"psrl.d") ||
346 Name.starts_with(
"psrl.q") ||
347 Name.starts_with(
"psrl.w") ||
348 Name.starts_with(
"psrli") ||
349 Name.starts_with(
"psrlv") ||
350 Name.starts_with(
"psub.") ||
351 Name.starts_with(
"psubs.") ||
352 Name.starts_with(
"psubus.") ||
353 Name.starts_with(
"pternlog.") ||
354 Name.starts_with(
"punpckh") ||
355 Name.starts_with(
"punpckl") ||
356 Name.starts_with(
"pxor.") ||
357 Name.starts_with(
"shuf.f") ||
358 Name.starts_with(
"shuf.i") ||
359 Name.starts_with(
"shuf.p") ||
360 Name.starts_with(
"sqrt.p") ||
361 Name.starts_with(
"store.b.") ||
362 Name.starts_with(
"store.d.") ||
363 Name.starts_with(
"store.p") ||
364 Name.starts_with(
"store.q.") ||
365 Name.starts_with(
"store.w.") ||
366 Name ==
"store.ss" ||
367 Name.starts_with(
"storeu.") ||
368 Name.starts_with(
"sub.p") ||
369 Name.starts_with(
"ucmp.") ||
370 Name.starts_with(
"unpckh.") ||
371 Name.starts_with(
"unpckl.") ||
372 Name.starts_with(
"valign.") ||
373 Name ==
"vcvtph2ps.128" ||
374 Name ==
"vcvtph2ps.256" ||
375 Name.starts_with(
"vextract") ||
376 Name.starts_with(
"vfmadd.") ||
377 Name.starts_with(
"vfmaddsub.") ||
378 Name.starts_with(
"vfnmadd.") ||
379 Name.starts_with(
"vfnmsub.") ||
380 Name.starts_with(
"vpdpbusd.") ||
381 Name.starts_with(
"vpdpbusds.") ||
382 Name.starts_with(
"vpdpwssd.") ||
383 Name.starts_with(
"vpdpwssds.") ||
384 Name.starts_with(
"vpermi2var.") ||
385 Name.starts_with(
"vpermil.p") ||
386 Name.starts_with(
"vpermilvar.") ||
387 Name.starts_with(
"vpermt2var.") ||
388 Name.starts_with(
"vpmadd52") ||
389 Name.starts_with(
"vpshld.") ||
390 Name.starts_with(
"vpshldv.") ||
391 Name.starts_with(
"vpshrd.") ||
392 Name.starts_with(
"vpshrdv.") ||
393 Name.starts_with(
"vpshufbitqmb.") ||
394 Name.starts_with(
"xor."));
396 if (Name.consume_front(
"mask3."))
398 return (Name.starts_with(
"vfmadd.") ||
399 Name.starts_with(
"vfmaddsub.") ||
400 Name.starts_with(
"vfmsub.") ||
401 Name.starts_with(
"vfmsubadd.") ||
402 Name.starts_with(
"vfnmsub."));
404 if (Name.consume_front(
"maskz."))
406 return (Name.starts_with(
"pternlog.") ||
407 Name.starts_with(
"vfmadd.") ||
408 Name.starts_with(
"vfmaddsub.") ||
409 Name.starts_with(
"vpdpbusd.") ||
410 Name.starts_with(
"vpdpbusds.") ||
411 Name.starts_with(
"vpdpwssd.") ||
412 Name.starts_with(
"vpdpwssds.") ||
413 Name.starts_with(
"vpermt2var.") ||
414 Name.starts_with(
"vpmadd52") ||
415 Name.starts_with(
"vpshldv.") ||
416 Name.starts_with(
"vpshrdv."));
419 return (Name ==
"movntdqa" ||
420 Name ==
"pmul.dq.512" ||
421 Name ==
"pmulu.dq.512" ||
422 Name.starts_with(
"broadcastm") ||
423 Name.starts_with(
"cmp.p") ||
424 Name.starts_with(
"cvtb2mask.") ||
425 Name.starts_with(
"cvtd2mask.") ||
426 Name.starts_with(
"cvtmask2") ||
427 Name.starts_with(
"cvtq2mask.") ||
428 Name ==
"cvtusi2sd" ||
429 Name.starts_with(
"cvtw2mask.") ||
434 Name ==
"kortestc.w" ||
435 Name ==
"kortestz.w" ||
436 Name.starts_with(
"kunpck") ||
439 Name.starts_with(
"padds.") ||
440 Name.starts_with(
"pbroadcast") ||
441 Name.starts_with(
"prol") ||
442 Name.starts_with(
"pror") ||
443 Name.starts_with(
"psll.dq") ||
444 Name.starts_with(
"psrl.dq") ||
445 Name.starts_with(
"psubs.") ||
446 Name.starts_with(
"ptestm") ||
447 Name.starts_with(
"ptestnm") ||
448 Name.starts_with(
"storent.") ||
449 Name.starts_with(
"vbroadcast.s") ||
450 Name.starts_with(
"vpshld.") ||
451 Name.starts_with(
"vpshrd."));
454 if (Name.consume_front(
"fma."))
455 return (Name.starts_with(
"vfmadd.") ||
456 Name.starts_with(
"vfmsub.") ||
457 Name.starts_with(
"vfmsubadd.") ||
458 Name.starts_with(
"vfnmadd.") ||
459 Name.starts_with(
"vfnmsub."));
461 if (Name.consume_front(
"fma4."))
462 return Name.starts_with(
"vfmadd.s");
464 if (Name.consume_front(
"sse."))
465 return (Name ==
"add.ss" ||
466 Name ==
"cvtsi2ss" ||
467 Name ==
"cvtsi642ss" ||
470 Name.starts_with(
"sqrt.p") ||
472 Name.starts_with(
"storeu.") ||
475 if (Name.consume_front(
"sse2."))
476 return (Name ==
"add.sd" ||
477 Name ==
"cvtdq2pd" ||
478 Name ==
"cvtdq2ps" ||
479 Name ==
"cvtps2pd" ||
480 Name ==
"cvtsi2sd" ||
481 Name ==
"cvtsi642sd" ||
482 Name ==
"cvtss2sd" ||
485 Name.starts_with(
"padds.") ||
486 Name.starts_with(
"paddus.") ||
487 Name.starts_with(
"pcmpeq.") ||
488 Name.starts_with(
"pcmpgt.") ||
493 Name ==
"pmulu.dq" ||
494 Name.starts_with(
"pshuf") ||
495 Name.starts_with(
"psll.dq") ||
496 Name.starts_with(
"psrl.dq") ||
497 Name.starts_with(
"psubs.") ||
498 Name.starts_with(
"psubus.") ||
499 Name.starts_with(
"sqrt.p") ||
501 Name ==
"storel.dq" ||
502 Name.starts_with(
"storeu.") ||
505 if (Name.consume_front(
"sse41."))
506 return (Name.starts_with(
"blendp") ||
507 Name ==
"movntdqa" ||
517 Name.starts_with(
"pmovsx") ||
518 Name.starts_with(
"pmovzx") ||
521 if (Name.consume_front(
"sse42."))
522 return Name ==
"crc32.64.8";
524 if (Name.consume_front(
"sse4a."))
525 return Name.starts_with(
"movnt.");
527 if (Name.consume_front(
"ssse3."))
528 return (Name ==
"pabs.b.128" ||
529 Name ==
"pabs.d.128" ||
530 Name ==
"pabs.w.128");
532 if (Name.consume_front(
"xop."))
533 return (Name ==
"vpcmov" ||
534 Name ==
"vpcmov.256" ||
535 Name.starts_with(
"vpcom") ||
536 Name.starts_with(
"vprot"));
538 if (Name.consume_front(
"bmi."))
539 return (Name.starts_with(
"pdep.") ||
540 Name.starts_with(
"pext."));
542 return (Name ==
"addcarry.u32" ||
543 Name ==
"addcarry.u64" ||
544 Name ==
"addcarryx.u32" ||
545 Name ==
"addcarryx.u64" ||
546 Name ==
"subborrow.u32" ||
547 Name ==
"subborrow.u64" ||
548 Name.starts_with(
"vcvtph2ps."));
554 if (!Name.consume_front(
"x86."))
562 if (Name ==
"rdtscp") {
564 if (
F->getFunctionType()->getNumParams() == 0)
569 Intrinsic::x86_rdtscp);
576 if (Name.consume_front(
"sse41.ptest")) {
578 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
579 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
580 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
593 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
594 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
595 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
596 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
597 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
598 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
603 if (Name.consume_front(
"avx512.")) {
604 if (Name.consume_front(
"mask.cmp.")) {
607 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
608 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
609 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
610 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
611 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
612 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
616 }
else if (Name.starts_with(
"vpdpbusd.") ||
617 Name.starts_with(
"vpdpbusds.")) {
620 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
621 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
622 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
623 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
624 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
625 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
629 }
else if (Name.starts_with(
"vpdpwssd.") ||
630 Name.starts_with(
"vpdpwssds.")) {
633 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
634 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
635 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
636 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
637 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
638 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
646 if (Name.consume_front(
"avx2.")) {
647 if (Name.consume_front(
"vpdpb")) {
650 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
651 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
652 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
653 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
654 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
655 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
656 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
657 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
658 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
659 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
660 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
661 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
665 }
else if (Name.consume_front(
"vpdpw")) {
668 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
669 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
670 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
671 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
672 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
673 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
674 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
675 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
676 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
677 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
678 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
679 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
687 if (Name.consume_front(
"avx10.")) {
688 if (Name.consume_front(
"vpdpb")) {
691 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
692 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
693 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
694 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
695 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
696 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
700 }
else if (Name.consume_front(
"vpdpw")) {
702 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
703 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
704 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
705 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
706 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
707 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
715 if (Name.consume_front(
"avx512bf16.")) {
718 .
Case(
"cvtne2ps2bf16.128",
719 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
720 .
Case(
"cvtne2ps2bf16.256",
721 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
722 .
Case(
"cvtne2ps2bf16.512",
723 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
724 .
Case(
"mask.cvtneps2bf16.128",
725 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
726 .
Case(
"cvtneps2bf16.256",
727 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
728 .
Case(
"cvtneps2bf16.512",
729 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
736 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
737 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
738 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
745 if (Name.consume_front(
"xop.")) {
747 if (Name.starts_with(
"vpermil2")) {
750 auto Idx =
F->getFunctionType()->getParamType(2);
751 if (Idx->isFPOrFPVectorTy()) {
752 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
753 unsigned EltSize = Idx->getScalarSizeInBits();
754 if (EltSize == 64 && IdxSize == 128)
755 ID = Intrinsic::x86_xop_vpermil2pd;
756 else if (EltSize == 32 && IdxSize == 128)
757 ID = Intrinsic::x86_xop_vpermil2ps;
758 else if (EltSize == 64 && IdxSize == 256)
759 ID = Intrinsic::x86_xop_vpermil2pd_256;
761 ID = Intrinsic::x86_xop_vpermil2ps_256;
763 }
else if (
F->arg_size() == 2)
766 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
767 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
778 if (Name ==
"seh.recoverfp") {
780 Intrinsic::eh_recoverfp);
792 if (Name.starts_with(
"rbit")) {
795 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
799 if (Name ==
"thread.pointer") {
802 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
806 bool Neon = Name.consume_front(
"neon.");
811 if (Name.consume_front(
"bfdot.")) {
815 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
820 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
821 assert((OperandWidth == 64 || OperandWidth == 128) &&
822 "Unexpected operand width");
824 std::array<Type *, 2> Tys{
835 if (Name.consume_front(
"bfm")) {
837 if (Name.consume_back(
".v4f32.v16i8")) {
883 F->arg_begin()->getType());
887 if (Name.consume_front(
"vst")) {
889 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
893 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
894 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
897 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
898 Intrinsic::arm_neon_vst4lane};
900 auto fArgs =
F->getFunctionType()->params();
901 Type *Tys[] = {fArgs[0], fArgs[1]};
904 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
907 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
916 if (Name.consume_front(
"mve.")) {
918 if (Name ==
"vctp64") {
928 if (Name.starts_with(
"vrintn.v")) {
930 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
935 if (Name.consume_back(
".v4i1")) {
937 if (Name.consume_back(
".predicated.v2i64.v4i32"))
939 return Name ==
"mull.int" || Name ==
"vqdmull";
941 if (Name.consume_back(
".v2i64")) {
943 bool IsGather = Name.consume_front(
"vldr.gather.");
944 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
945 if (Name.consume_front(
"base.")) {
947 Name.consume_front(
"wb.");
950 return Name ==
"predicated.v2i64";
953 if (Name.consume_front(
"offset.predicated."))
954 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
955 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
968 if (Name.consume_front(
"cde.vcx")) {
970 if (Name.consume_back(
".predicated.v2i64.v4i1"))
972 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
973 Name ==
"3q" || Name ==
"3qa";
987 F->arg_begin()->getType());
991 if (Name.starts_with(
"addp")) {
993 if (
F->arg_size() != 2)
996 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
998 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1004 if (Name.starts_with(
"bfcvt")) {
1010 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1017 if (Name.consume_front(
"sve.")) {
1019 if (Name.consume_front(
"bf")) {
1020 if (Name ==
"mmla") {
1021 Type *Tys[] = {
F->getReturnType(),
1022 std::next(
F->arg_begin())->getType()};
1024 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1027 if (Name.consume_back(
".lane")) {
1031 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1032 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1033 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1045 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1050 if (Name.consume_front(
"convert.from.svbool")) {
1053 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1056 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1061 if (Name.consume_front(
"convert.to.svbool")) {
1064 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1067 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1072 if (Name.consume_front(
"addqv")) {
1074 if (!
F->getReturnType()->isFPOrFPVectorTy())
1077 auto Args =
F->getFunctionType()->params();
1078 Type *Tys[] = {
F->getReturnType(), Args[1]};
1080 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1084 if (Name.consume_front(
"ld")) {
1086 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1087 if (LdRegex.
match(Name)) {
1093 "Expected 2 arguments for ld* intrinsic.");
1094 Type *PtrTy =
F->getArg(1)->getType();
1097 Intrinsic::aarch64_sve_ld2_sret,
1098 Intrinsic::aarch64_sve_ld3_sret,
1099 Intrinsic::aarch64_sve_ld4_sret,
1102 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1108 if (Name.consume_front(
"tuple.")) {
1110 if (Name.starts_with(
"get")) {
1112 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1114 F->getParent(), Intrinsic::vector_extract, Tys);
1118 if (Name.starts_with(
"set")) {
1120 auto Args =
F->getFunctionType()->params();
1121 Type *Tys[] = {Args[0], Args[2], Args[1]};
1123 F->getParent(), Intrinsic::vector_insert, Tys);
1127 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1128 if (CreateTupleRegex.
match(Name)) {
1130 auto Args =
F->getFunctionType()->params();
1131 Type *Tys[] = {
F->getReturnType(), Args[1]};
1133 F->getParent(), Intrinsic::vector_insert, Tys);
1139 if (Name.starts_with(
"rev.nxv")) {
1142 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1148 if (Name.consume_front(
"sme.")) {
1150 if (Name.consume_front(
"ftmopa.")) {
1155 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1156 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1173 if (Name.consume_front(
"cp.async.bulk.tensor.g2s.")) {
1177 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d)
1179 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d)
1181 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d)
1182 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d)
1183 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d)
1184 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d)
1185 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d)
1186 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d)
1195 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1209 size_t FlagStartIndex =
F->getFunctionType()->getNumParams() - 3;
1210 Type *ArgType =
F->getFunctionType()->getParamType(FlagStartIndex);
1235 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1238 auto [RedOpName, ShapeName] = Name.split(
'.');
1243 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1244 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1245 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1246 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1247 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1248 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1249 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1250 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1256 if (Name.consume_front(
"mapa.shared.cluster"))
1257 if (
F->getReturnType()->getPointerAddressSpace() ==
1259 return Intrinsic::nvvm_mapa_shared_cluster;
1261 if (Name.consume_front(
"cp.async.bulk.")) {
1264 .
Case(
"global.to.shared.cluster",
1265 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1266 .
Case(
"shared.cta.to.cluster",
1267 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1271 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1281 if (!Name.consume_front(
"tcgen05.commit."))
1284 if (Name.consume_front(
"shared."))
1286 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1287 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1290 if (Name.consume_front(
"mc.shared.")) {
1292 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1296 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1297 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1306 if (
F->arg_size() != 2)
1309 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1310 Name.consume_front(
"tcgen05.alloc."))
1312 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1313 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1316 if (Name.consume_front(
"tcgen05.dealloc."))
1318 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1319 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1326 if (Name.consume_front(
"fma.rn."))
1328 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1329 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1330 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1331 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1334 if (Name.consume_front(
"fmax."))
1336 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1337 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1338 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1339 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1340 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1341 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1342 .
Case(
"ftz.nan.xorsign.abs.bf16",
1343 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1344 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1345 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1346 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1347 .
Case(
"ftz.xorsign.abs.bf16x2",
1348 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1349 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1350 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1351 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1352 .
Case(
"nan.xorsign.abs.bf16x2",
1353 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1354 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1355 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1358 if (Name.consume_front(
"fmin."))
1360 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1361 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1362 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1363 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1364 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1365 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1366 .
Case(
"ftz.nan.xorsign.abs.bf16",
1367 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1368 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1369 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1370 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1371 .
Case(
"ftz.xorsign.abs.bf16x2",
1372 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1373 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1374 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1375 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1376 .
Case(
"nan.xorsign.abs.bf16x2",
1377 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1378 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1379 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1382 if (Name.consume_front(
"neg."))
1384 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1385 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1393 if (!Name.consume_front(
"tcgen05.mma."))
1397 if (Name.starts_with(
"ws"))
1400 return F->getIntrinsicID();
1404 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1405 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1406 Name.consume_front(
"param");
1410 if (!Name.consume_front(
"vp."))
1439 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1440 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1447 if (!Name.consume_front(
"vp."))
1467 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1468 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1473 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1485 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1496 if (Name.starts_with(
"to.fp16")) {
1500 FuncTy->getReturnType());
1503 if (Name.starts_with(
"from.fp16")) {
1507 FuncTy->getReturnType());
1519 if (Defaults.empty())
1531 if (
F->arg_size() >= FullDecl->
arg_size())
1536 if (
F->arg_size() < FirstDefault)
1544 bool CanUpgradeDebugIntrinsicsToRecords) {
1545 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1550 if (!Name.consume_front(
"llvm.") || Name.empty())
1556 bool IsArm = Name.consume_front(
"arm.");
1557 if (IsArm || Name.consume_front(
"aarch64.")) {
1563 if (Name.consume_front(
"amdgcn.")) {
1564 if (Name ==
"alignbit") {
1567 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1571 if (Name.consume_front(
"atomic.")) {
1572 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1573 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1582 switch (
F->getIntrinsicID()) {
1586 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1587 if (
F->arg_size() == 7) {
1592 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1593 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1594 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1595 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1596 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1597 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1598 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1599 if (
F->arg_size() == 8) {
1606 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1607 Name.consume_front(
"flat.atomic.")) {
1608 if (Name.starts_with(
"fadd") ||
1610 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1611 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1619 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1624 if (Name.starts_with(
"ldexp.")) {
1627 F->getParent(), Intrinsic::ldexp,
1628 {F->getReturnType(), F->getArg(1)->getType()});
1637 if (
F->arg_size() == 1) {
1638 if (Name.consume_front(
"convert.")) {
1652 F->arg_begin()->getType());
1658 if (Name ==
"coro.end" &&
1659 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1660 CoroEndID = Intrinsic::coro_end;
1661 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1662 CoroEndID = Intrinsic::coro_end_async;
1673 if (Name.consume_front(
"dbg.")) {
1675 if (CanUpgradeDebugIntrinsicsToRecords) {
1676 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1677 Name ==
"declare" || Name ==
"label") {
1686 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1689 Intrinsic::dbg_value);
1696 if (Name.consume_front(
"experimental.vector.")) {
1702 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1703 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1704 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1705 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1706 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1708 Intrinsic::vector_partial_reduce_add)
1711 const auto *FT =
F->getFunctionType();
1713 if (ID == Intrinsic::vector_extract ||
1714 ID == Intrinsic::vector_interleave2)
1717 if (ID != Intrinsic::vector_interleave2)
1719 if (ID == Intrinsic::vector_insert ||
1720 ID == Intrinsic::vector_partial_reduce_add)
1728 if (Name.consume_front(
"reduce.")) {
1730 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1731 if (R.match(Name, &
Groups))
1733 .
Case(
"add", Intrinsic::vector_reduce_add)
1734 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1735 .
Case(
"and", Intrinsic::vector_reduce_and)
1736 .
Case(
"or", Intrinsic::vector_reduce_or)
1737 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1738 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1739 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1740 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1741 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1742 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1743 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1748 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1753 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1754 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1759 auto Args =
F->getFunctionType()->params();
1761 {Args[V2 ? 1 : 0]});
1767 if (Name.consume_front(
"splice"))
1771 if (Name.consume_front(
"experimental.stepvector.")) {
1775 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1780 if (Name.starts_with(
"flt.rounds")) {
1783 Intrinsic::get_rounding);
1788 if (Name.starts_with(
"invariant.group.barrier")) {
1790 auto Args =
F->getFunctionType()->params();
1791 Type* ObjectPtr[1] = {Args[0]};
1794 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1799 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1800 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1801 if (IsLifetimeStart || IsLifetimeEnd) {
1802 if (
F->arg_size() == 2) {
1803 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1804 : Intrinsic::lifetime_end;
1809 F->getArg(1)->getType());
1811 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1831 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1832 .StartsWith(
"memmove.", Intrinsic::memmove)
1834 if (
F->arg_size() == 5) {
1838 F->getFunctionType()->params().slice(0, 3);
1844 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1847 const auto *FT =
F->getFunctionType();
1848 Type *ParamTypes[2] = {
1849 FT->getParamType(0),
1853 Intrinsic::memset, ParamTypes);
1859 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1860 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1861 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1862 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1864 if (MaskedID &&
F->arg_size() == 4) {
1866 if (MaskedID == Intrinsic::masked_load ||
1867 MaskedID == Intrinsic::masked_gather) {
1869 F->getParent(), MaskedID,
1870 {F->getReturnType(), F->getArg(0)->getType()});
1874 F->getParent(), MaskedID,
1875 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1881 if (Name.consume_front(
"nvvm.")) {
1883 if (
F->arg_size() == 1) {
1886 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1887 .Case(
"clz.i", Intrinsic::ctlz)
1888 .
Case(
"popc.i", Intrinsic::ctpop)
1892 {F->getReturnType()});
1895 }
else if (
F->arg_size() == 2) {
1898 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1899 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1900 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1901 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1905 {F->getReturnType()});
1911 if (!
F->getReturnType()->getScalarType()->isBFloatTy()) {
1941 F->getParent(), IID,
F->getReturnType(),
1942 F->getFunctionType()->params());
1953 {F->getArg(0)->getType()});
1978 bool Expand =
false;
1979 if (Name.consume_front(
"abs."))
1982 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
1983 else if (Name.consume_front(
"fabs."))
1985 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
1986 else if (Name.consume_front(
"ex2.approx."))
1989 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
1990 else if (Name.consume_front(
"atomic.load."))
1999 else if (Name.consume_front(
"atomic."))
2014 else if (Name.consume_front(
"bitcast."))
2017 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2018 else if (Name.consume_front(
"rotate."))
2020 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2021 else if (Name.consume_front(
"ptr.gen.to."))
2024 else if (Name.consume_front(
"ptr."))
2027 else if (Name.consume_front(
"ldg.global."))
2029 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2030 Name.starts_with(
"p."));
2033 .
Case(
"barrier0",
true)
2034 .
Case(
"barrier.n",
true)
2035 .
Case(
"barrier.sync.cnt",
true)
2036 .
Case(
"barrier.sync",
true)
2037 .
Case(
"barrier",
true)
2038 .
Case(
"bar.sync",
true)
2039 .
Case(
"barrier0.popc",
true)
2040 .
Case(
"barrier0.and",
true)
2041 .
Case(
"barrier0.or",
true)
2042 .
Case(
"clz.ll",
true)
2043 .
Case(
"popc.ll",
true)
2045 .
Case(
"swap.lo.hi.b64",
true)
2046 .
Case(
"tanh.approx.f32",
true)
2058 if (Name.starts_with(
"objectsize.")) {
2059 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2060 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2063 Intrinsic::objectsize, Tys);
2070 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2073 F->getParent(), Intrinsic::ptr_annotation,
2074 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2080 if (Name.consume_front(
"riscv.")) {
2083 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2084 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2085 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2086 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2089 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2102 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2103 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2112 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2113 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2114 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2115 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2120 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2129 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2131 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2140 if (Name ==
"stackprotectorcheck") {
2147 if (Name ==
"thread.pointer") {
2149 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2155 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2158 F->getParent(), Intrinsic::var_annotation,
2159 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2162 if (Name.consume_front(
"vector.splice")) {
2163 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2173 if (Name.consume_front(
"wasm.")) {
2176 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2177 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2178 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2183 F->getReturnType());
2187 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2189 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2191 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2210 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2220 std::string
Name =
F->getName().str();
2223 Name,
F->getParent());
2234 if (Result != std::nullopt) {
2250 bool CanUpgradeDebugIntrinsicsToRecords) {
2270 GV->
getName() ==
"llvm.global_dtors")) ||
2285 unsigned N =
Init->getNumOperands();
2286 std::vector<Constant *> NewCtors(
N);
2287 for (
unsigned i = 0; i !=
N; ++i) {
2290 Ctor->getAggregateElement(1),
2304 unsigned NumElts = ResultTy->getNumElements() * 8;
2308 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2318 for (
unsigned l = 0; l != NumElts; l += 16)
2319 for (
unsigned i = 0; i != 16; ++i) {
2320 unsigned Idx = NumElts + i - Shift;
2322 Idx -= NumElts - 16;
2323 Idxs[l + i] = Idx + l;
2326 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2330 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2338 unsigned NumElts = ResultTy->getNumElements() * 8;
2342 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2352 for (
unsigned l = 0; l != NumElts; l += 16)
2353 for (
unsigned i = 0; i != 16; ++i) {
2354 unsigned Idx = i + Shift;
2356 Idx += NumElts - 16;
2357 Idxs[l + i] = Idx + l;
2360 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2364 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2372 Mask = Builder.CreateBitCast(Mask, MaskTy);
2378 for (
unsigned i = 0; i != NumElts; ++i)
2380 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2391 if (
C->isAllOnesValue())
2396 return Builder.CreateSelect(Mask, Op0, Op1);
2403 if (
C->isAllOnesValue())
2407 Mask->getType()->getIntegerBitWidth());
2408 Mask = Builder.CreateBitCast(Mask, MaskTy);
2409 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2410 return Builder.CreateSelect(Mask, Op0, Op1);
2423 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2424 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2429 ShiftVal &= (NumElts - 1);
2438 if (ShiftVal > 16) {
2446 for (
unsigned l = 0; l < NumElts; l += 16) {
2447 for (
unsigned i = 0; i != 16; ++i) {
2448 unsigned Idx = ShiftVal + i;
2449 if (!IsVALIGN && Idx >= 16)
2450 Idx += NumElts - 16;
2451 Indices[l + i] = Idx + l;
2456 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2462 bool ZeroMask,
bool IndexForm) {
2465 unsigned EltWidth = Ty->getScalarSizeInBits();
2466 bool IsFloat = Ty->isFPOrFPVectorTy();
2468 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2469 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2470 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2471 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2472 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2473 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2474 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2475 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2476 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2477 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2478 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2479 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2480 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2481 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2482 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2483 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2484 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2485 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2486 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2487 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2488 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2489 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2490 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2491 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2492 else if (VecWidth == 128 && EltWidth == 16)
2493 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2494 else if (VecWidth == 256 && EltWidth == 16)
2495 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2496 else if (VecWidth == 512 && EltWidth == 16)
2497 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2498 else if (VecWidth == 128 && EltWidth == 8)
2499 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2500 else if (VecWidth == 256 && EltWidth == 8)
2501 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2502 else if (VecWidth == 512 && EltWidth == 8)
2503 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2514 Value *V = Builder.CreateIntrinsic(IID, Args);
2526 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2537 bool IsRotateRight) {
2547 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2548 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2551 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2552 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2597 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2602 bool IsShiftRight,
bool ZeroMask) {
2616 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2617 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2620 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2621 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2636 const Align Alignment =
2638 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2643 if (
C->isAllOnesValue())
2644 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2649 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2655 const Align Alignment =
2664 if (
C->isAllOnesValue())
2665 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2670 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2676 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2677 {Op0, Builder.getInt1(
false)});
2692 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2693 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2694 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2695 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2696 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2699 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2700 LHS = Builder.CreateAnd(
LHS, Mask);
2701 RHS = Builder.CreateAnd(
RHS, Mask);
2718 if (!
C || !
C->isAllOnesValue())
2719 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2724 for (
unsigned i = 0; i != NumElts; ++i)
2726 for (
unsigned i = NumElts; i != 8; ++i)
2727 Indices[i] = NumElts + i % NumElts;
2728 Vec = Builder.CreateShuffleVector(Vec,
2732 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2736 unsigned CC,
bool Signed) {
2744 }
else if (CC == 7) {
2780 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2781 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2783 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2784 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2793 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2799 Name = Name.substr(12);
2804 if (Name.starts_with(
"max.p")) {
2805 if (VecWidth == 128 && EltWidth == 32)
2806 IID = Intrinsic::x86_sse_max_ps;
2807 else if (VecWidth == 128 && EltWidth == 64)
2808 IID = Intrinsic::x86_sse2_max_pd;
2809 else if (VecWidth == 256 && EltWidth == 32)
2810 IID = Intrinsic::x86_avx_max_ps_256;
2811 else if (VecWidth == 256 && EltWidth == 64)
2812 IID = Intrinsic::x86_avx_max_pd_256;
2815 }
else if (Name.starts_with(
"min.p")) {
2816 if (VecWidth == 128 && EltWidth == 32)
2817 IID = Intrinsic::x86_sse_min_ps;
2818 else if (VecWidth == 128 && EltWidth == 64)
2819 IID = Intrinsic::x86_sse2_min_pd;
2820 else if (VecWidth == 256 && EltWidth == 32)
2821 IID = Intrinsic::x86_avx_min_ps_256;
2822 else if (VecWidth == 256 && EltWidth == 64)
2823 IID = Intrinsic::x86_avx_min_pd_256;
2826 }
else if (Name.starts_with(
"pshuf.b.")) {
2827 if (VecWidth == 128)
2828 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2829 else if (VecWidth == 256)
2830 IID = Intrinsic::x86_avx2_pshuf_b;
2831 else if (VecWidth == 512)
2832 IID = Intrinsic::x86_avx512_pshuf_b_512;
2835 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2836 if (VecWidth == 128)
2837 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2838 else if (VecWidth == 256)
2839 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2840 else if (VecWidth == 512)
2841 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2844 }
else if (Name.starts_with(
"pmulh.w.")) {
2845 if (VecWidth == 128)
2846 IID = Intrinsic::x86_sse2_pmulh_w;
2847 else if (VecWidth == 256)
2848 IID = Intrinsic::x86_avx2_pmulh_w;
2849 else if (VecWidth == 512)
2850 IID = Intrinsic::x86_avx512_pmulh_w_512;
2853 }
else if (Name.starts_with(
"pmulhu.w.")) {
2854 if (VecWidth == 128)
2855 IID = Intrinsic::x86_sse2_pmulhu_w;
2856 else if (VecWidth == 256)
2857 IID = Intrinsic::x86_avx2_pmulhu_w;
2858 else if (VecWidth == 512)
2859 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2862 }
else if (Name.starts_with(
"pmaddw.d.")) {
2863 if (VecWidth == 128)
2864 IID = Intrinsic::x86_sse2_pmadd_wd;
2865 else if (VecWidth == 256)
2866 IID = Intrinsic::x86_avx2_pmadd_wd;
2867 else if (VecWidth == 512)
2868 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2871 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2872 if (VecWidth == 128)
2873 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2874 else if (VecWidth == 256)
2875 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2876 else if (VecWidth == 512)
2877 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2880 }
else if (Name.starts_with(
"packsswb.")) {
2881 if (VecWidth == 128)
2882 IID = Intrinsic::x86_sse2_packsswb_128;
2883 else if (VecWidth == 256)
2884 IID = Intrinsic::x86_avx2_packsswb;
2885 else if (VecWidth == 512)
2886 IID = Intrinsic::x86_avx512_packsswb_512;
2889 }
else if (Name.starts_with(
"packssdw.")) {
2890 if (VecWidth == 128)
2891 IID = Intrinsic::x86_sse2_packssdw_128;
2892 else if (VecWidth == 256)
2893 IID = Intrinsic::x86_avx2_packssdw;
2894 else if (VecWidth == 512)
2895 IID = Intrinsic::x86_avx512_packssdw_512;
2898 }
else if (Name.starts_with(
"packuswb.")) {
2899 if (VecWidth == 128)
2900 IID = Intrinsic::x86_sse2_packuswb_128;
2901 else if (VecWidth == 256)
2902 IID = Intrinsic::x86_avx2_packuswb;
2903 else if (VecWidth == 512)
2904 IID = Intrinsic::x86_avx512_packuswb_512;
2907 }
else if (Name.starts_with(
"packusdw.")) {
2908 if (VecWidth == 128)
2909 IID = Intrinsic::x86_sse41_packusdw;
2910 else if (VecWidth == 256)
2911 IID = Intrinsic::x86_avx2_packusdw;
2912 else if (VecWidth == 512)
2913 IID = Intrinsic::x86_avx512_packusdw_512;
2916 }
else if (Name.starts_with(
"vpermilvar.")) {
2917 if (VecWidth == 128 && EltWidth == 32)
2918 IID = Intrinsic::x86_avx_vpermilvar_ps;
2919 else if (VecWidth == 128 && EltWidth == 64)
2920 IID = Intrinsic::x86_avx_vpermilvar_pd;
2921 else if (VecWidth == 256 && EltWidth == 32)
2922 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2923 else if (VecWidth == 256 && EltWidth == 64)
2924 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2925 else if (VecWidth == 512 && EltWidth == 32)
2926 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2927 else if (VecWidth == 512 && EltWidth == 64)
2928 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2931 }
else if (Name ==
"cvtpd2dq.256") {
2932 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2933 }
else if (Name ==
"cvtpd2ps.256") {
2934 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2935 }
else if (Name ==
"cvttpd2dq.256") {
2936 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2937 }
else if (Name ==
"cvttps2dq.128") {
2938 IID = Intrinsic::x86_sse2_cvttps2dq;
2939 }
else if (Name ==
"cvttps2dq.256") {
2940 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2941 }
else if (Name.starts_with(
"permvar.")) {
2943 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2944 IID = Intrinsic::x86_avx2_permps;
2945 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2946 IID = Intrinsic::x86_avx2_permd;
2947 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2948 IID = Intrinsic::x86_avx512_permvar_df_256;
2949 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2950 IID = Intrinsic::x86_avx512_permvar_di_256;
2951 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2952 IID = Intrinsic::x86_avx512_permvar_sf_512;
2953 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2954 IID = Intrinsic::x86_avx512_permvar_si_512;
2955 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2956 IID = Intrinsic::x86_avx512_permvar_df_512;
2957 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2958 IID = Intrinsic::x86_avx512_permvar_di_512;
2959 else if (VecWidth == 128 && EltWidth == 16)
2960 IID = Intrinsic::x86_avx512_permvar_hi_128;
2961 else if (VecWidth == 256 && EltWidth == 16)
2962 IID = Intrinsic::x86_avx512_permvar_hi_256;
2963 else if (VecWidth == 512 && EltWidth == 16)
2964 IID = Intrinsic::x86_avx512_permvar_hi_512;
2965 else if (VecWidth == 128 && EltWidth == 8)
2966 IID = Intrinsic::x86_avx512_permvar_qi_128;
2967 else if (VecWidth == 256 && EltWidth == 8)
2968 IID = Intrinsic::x86_avx512_permvar_qi_256;
2969 else if (VecWidth == 512 && EltWidth == 8)
2970 IID = Intrinsic::x86_avx512_permvar_qi_512;
2973 }
else if (Name.starts_with(
"dbpsadbw.")) {
2974 if (VecWidth == 128)
2975 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2976 else if (VecWidth == 256)
2977 IID = Intrinsic::x86_avx512_dbpsadbw_256;
2978 else if (VecWidth == 512)
2979 IID = Intrinsic::x86_avx512_dbpsadbw_512;
2982 }
else if (Name.starts_with(
"pmultishift.qb.")) {
2983 if (VecWidth == 128)
2984 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
2985 else if (VecWidth == 256)
2986 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
2987 else if (VecWidth == 512)
2988 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
2991 }
else if (Name.starts_with(
"conflict.")) {
2992 if (Name[9] ==
'd' && VecWidth == 128)
2993 IID = Intrinsic::x86_avx512_conflict_d_128;
2994 else if (Name[9] ==
'd' && VecWidth == 256)
2995 IID = Intrinsic::x86_avx512_conflict_d_256;
2996 else if (Name[9] ==
'd' && VecWidth == 512)
2997 IID = Intrinsic::x86_avx512_conflict_d_512;
2998 else if (Name[9] ==
'q' && VecWidth == 128)
2999 IID = Intrinsic::x86_avx512_conflict_q_128;
3000 else if (Name[9] ==
'q' && VecWidth == 256)
3001 IID = Intrinsic::x86_avx512_conflict_q_256;
3002 else if (Name[9] ==
'q' && VecWidth == 512)
3003 IID = Intrinsic::x86_avx512_conflict_q_512;
3006 }
else if (Name.starts_with(
"pavg.")) {
3007 if (Name[5] ==
'b' && VecWidth == 128)
3008 IID = Intrinsic::x86_sse2_pavg_b;
3009 else if (Name[5] ==
'b' && VecWidth == 256)
3010 IID = Intrinsic::x86_avx2_pavg_b;
3011 else if (Name[5] ==
'b' && VecWidth == 512)
3012 IID = Intrinsic::x86_avx512_pavg_b_512;
3013 else if (Name[5] ==
'w' && VecWidth == 128)
3014 IID = Intrinsic::x86_sse2_pavg_w;
3015 else if (Name[5] ==
'w' && VecWidth == 256)
3016 IID = Intrinsic::x86_avx2_pavg_w;
3017 else if (Name[5] ==
'w' && VecWidth == 512)
3018 IID = Intrinsic::x86_avx512_pavg_w_512;
3027 Rep = Builder.CreateIntrinsic(IID, Args);
3038 if (AsmStr->find(
"mov\tfp") == 0 &&
3039 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3040 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3041 AsmStr->replace(Pos, 1,
";");
3047 Value *Rep =
nullptr;
3049 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3051 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3052 {Arg, Builder.getTrue()},
3054 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3055 Type *Ty = (Name ==
"abs.bf16")
3059 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3060 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3061 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3062 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3063 : Intrinsic::nvvm_fabs;
3064 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3065 }
else if (Name.consume_front(
"ex2.approx.")) {
3067 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3068 : Intrinsic::nvvm_ex2_approx;
3069 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3070 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3071 Name.starts_with(
"atomic.load.add.f64.p")) {
3074 Rep = Builder.CreateAtomicRMW(
3080 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3081 Name.starts_with(
"atomic.load.dec.32.p")) {
3086 Rep = Builder.CreateAtomicRMW(
3090 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3096 Op.contains(
".cta.") ?
"block" :
"");
3097 if (
Op.starts_with(
"cas.")) {
3099 Value *Pair = Builder.CreateAtomicCmpXchg(
3102 Rep = Builder.CreateExtractValue(Pair, 0);
3120 "unexpected nvvm scoped atomic intrinsic");
3121 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3124 }
else if (Name ==
"clz.ll") {
3127 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3128 {Arg, Builder.getFalse()},
3130 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3131 }
else if (Name ==
"popc.ll") {
3135 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3136 Arg,
nullptr,
"ctpop");
3137 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3138 }
else if (Name ==
"h2f") {
3140 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3141 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3142 }
else if (Name.consume_front(
"bitcast.") &&
3143 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3146 }
else if (Name ==
"rotate.b32") {
3149 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3150 {Arg, Arg, ShiftAmt});
3151 }
else if (Name ==
"rotate.b64") {
3155 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3156 {Arg, Arg, ZExtShiftAmt});
3157 }
else if (Name ==
"rotate.right.b64") {
3161 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3162 {Arg, Arg, ZExtShiftAmt});
3163 }
else if (Name ==
"swap.lo.hi.b64") {
3166 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3167 {Arg, Arg, Builder.getInt64(32)});
3168 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3171 Name.starts_with(
".to.gen"))) {
3173 }
else if (Name.consume_front(
"ldg.global")) {
3177 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3180 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3182 }
else if (Name ==
"tanh.approx.f32") {
3186 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3188 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3190 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3191 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3193 }
else if (Name ==
"barrier") {
3194 Rep = Builder.CreateIntrinsic(
3195 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3197 }
else if (Name ==
"barrier.sync") {
3198 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3200 }
else if (Name ==
"barrier.sync.cnt") {
3201 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3203 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3204 Name ==
"barrier0.or") {
3206 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3210 .
Case(
"barrier0.popc",
3211 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3212 .
Case(
"barrier0.and",
3213 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3214 .
Case(
"barrier0.or",
3215 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3216 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3217 Rep = Builder.CreateZExt(Bar, CI->
getType());
3221 !
F->getReturnType()->getScalarType()->isBFloatTy()) {
3231 ? Builder.CreateBitCast(Arg, NewType)
3234 Rep = Builder.CreateCall(NewFn, Args);
3235 if (
F->getReturnType()->isIntegerTy())
3236 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3246 Value *Rep =
nullptr;
3248 if (Name.starts_with(
"sse4a.movnt.")) {
3260 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3263 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3264 }
else if (Name.starts_with(
"avx.movnt.") ||
3265 Name.starts_with(
"avx512.storent.")) {
3277 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3278 }
else if (Name ==
"sse2.storel.dq") {
3283 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3284 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3285 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3286 }
else if (Name.starts_with(
"sse.storeu.") ||
3287 Name.starts_with(
"sse2.storeu.") ||
3288 Name.starts_with(
"avx.storeu.")) {
3291 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3292 }
else if (Name ==
"avx512.mask.store.ss") {
3296 }
else if (Name.starts_with(
"avx512.mask.store")) {
3298 bool Aligned = Name[17] !=
'u';
3301 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3304 bool CmpEq = Name[9] ==
'e';
3307 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3308 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3315 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3316 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3318 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3319 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3320 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3321 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3322 Name.starts_with(
"sse2.sqrt.p") ||
3323 Name.starts_with(
"sse.sqrt.p")) {
3324 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3325 {CI->getArgOperand(0)});
3326 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3330 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3331 : Intrinsic::x86_avx512_sqrt_pd_512;
3334 Rep = Builder.CreateIntrinsic(IID, Args);
3336 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3337 {CI->getArgOperand(0)});
3341 }
else if (Name.starts_with(
"avx512.ptestm") ||
3342 Name.starts_with(
"avx512.ptestnm")) {
3346 Rep = Builder.CreateAnd(Op0, Op1);
3352 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3354 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3357 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3360 }
else if (Name.starts_with(
"avx512.kunpck")) {
3365 for (
unsigned i = 0; i != NumElts; ++i)
3374 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3375 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3376 }
else if (Name ==
"avx512.kand.w") {
3379 Rep = Builder.CreateAnd(
LHS,
RHS);
3380 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3381 }
else if (Name ==
"avx512.kandn.w") {
3384 LHS = Builder.CreateNot(
LHS);
3385 Rep = Builder.CreateAnd(
LHS,
RHS);
3386 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3387 }
else if (Name ==
"avx512.kor.w") {
3390 Rep = Builder.CreateOr(
LHS,
RHS);
3391 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3392 }
else if (Name ==
"avx512.kxor.w") {
3395 Rep = Builder.CreateXor(
LHS,
RHS);
3396 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3397 }
else if (Name ==
"avx512.kxnor.w") {
3400 LHS = Builder.CreateNot(
LHS);
3401 Rep = Builder.CreateXor(
LHS,
RHS);
3402 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3403 }
else if (Name ==
"avx512.knot.w") {
3405 Rep = Builder.CreateNot(Rep);
3406 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3407 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3410 Rep = Builder.CreateOr(
LHS,
RHS);
3411 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3413 if (Name[14] ==
'c')
3417 Rep = Builder.CreateICmpEQ(Rep,
C);
3418 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3419 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3420 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3421 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3422 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3425 ConstantInt::get(I32Ty, 0));
3427 ConstantInt::get(I32Ty, 0));
3429 if (Name.contains(
".add."))
3430 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3431 else if (Name.contains(
".sub."))
3432 EltOp = Builder.CreateFSub(Elt0, Elt1);
3433 else if (Name.contains(
".mul."))
3434 EltOp = Builder.CreateFMul(Elt0, Elt1);
3436 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3437 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3438 ConstantInt::get(I32Ty, 0));
3439 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3441 bool CmpEq = Name[16] ==
'e';
3443 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3445 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3452 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3455 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3458 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3465 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3467 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3468 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3470 if (VecWidth == 128 && EltWidth == 32)
3471 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3472 else if (VecWidth == 256 && EltWidth == 32)
3473 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3474 else if (VecWidth == 512 && EltWidth == 32)
3475 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3476 else if (VecWidth == 128 && EltWidth == 64)
3477 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3478 else if (VecWidth == 256 && EltWidth == 64)
3479 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3480 else if (VecWidth == 512 && EltWidth == 64)
3481 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3488 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3491 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3492 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3494 if (VecWidth == 128 && EltWidth == 32)
3495 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3496 else if (VecWidth == 256 && EltWidth == 32)
3497 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3498 else if (VecWidth == 512 && EltWidth == 32)
3499 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3500 else if (VecWidth == 128 && EltWidth == 64)
3501 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3502 else if (VecWidth == 256 && EltWidth == 64)
3503 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3504 else if (VecWidth == 512 && EltWidth == 64)
3505 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3510 if (VecWidth == 512)
3512 Args.push_back(Mask);
3514 Rep = Builder.CreateIntrinsic(IID, Args);
3515 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3519 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3522 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3523 Name.starts_with(
"avx512.cvtw2mask.") ||
3524 Name.starts_with(
"avx512.cvtd2mask.") ||
3525 Name.starts_with(
"avx512.cvtq2mask.")) {
3530 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3531 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3532 Name.starts_with(
"avx512.mask.pabs")) {
3534 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3535 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3536 Name.starts_with(
"avx512.mask.pmaxs")) {
3538 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3539 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3540 Name.starts_with(
"avx512.mask.pmaxu")) {
3542 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3543 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3544 Name.starts_with(
"avx512.mask.pmins")) {
3546 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3547 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3548 Name.starts_with(
"avx512.mask.pminu")) {
3550 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3551 Name ==
"avx512.pmulu.dq.512" ||
3552 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3554 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3555 Name ==
"avx512.pmul.dq.512" ||
3556 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3558 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3559 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3564 }
else if (Name ==
"avx512.cvtusi2sd") {
3569 }
else if (Name ==
"sse2.cvtss2sd") {
3571 Rep = Builder.CreateFPExt(
3574 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3575 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3576 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3577 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3578 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3579 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3580 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3581 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3582 Name ==
"avx512.mask.cvtqq2ps.256" ||
3583 Name ==
"avx512.mask.cvtqq2ps.512" ||
3584 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3585 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3586 Name ==
"avx.cvt.ps2.pd.256" ||
3587 Name ==
"avx512.mask.cvtps2pd.128" ||
3588 Name ==
"avx512.mask.cvtps2pd.256") {
3593 unsigned NumDstElts = DstTy->getNumElements();
3594 if (NumDstElts < SrcTy->getNumElements()) {
3595 assert(NumDstElts == 2 &&
"Unexpected vector size");
3596 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3599 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3600 bool IsUnsigned = Name.contains(
"cvtu");
3602 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3606 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3607 : Intrinsic::x86_avx512_sitofp_round;
3608 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3611 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3612 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3618 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3619 Name.starts_with(
"vcvtph2ps.")) {
3623 unsigned NumDstElts = DstTy->getNumElements();
3624 if (NumDstElts != SrcTy->getNumElements()) {
3625 assert(NumDstElts == 4 &&
"Unexpected vector size");
3626 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3628 Rep = Builder.CreateBitCast(
3630 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3634 }
else if (Name.starts_with(
"avx512.mask.load")) {
3636 bool Aligned = Name[16] !=
'u';
3639 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3643 ResultTy->getNumElements());
3644 Rep = Builder.CreateIntrinsic(
3645 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3647 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3653 Rep = Builder.CreateIntrinsic(
3654 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3656 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3657 Name.starts_with(
"avx512.mask.expand.")) {
3661 ResultTy->getNumElements());
3663 bool IsCompress = Name[12] ==
'c';
3664 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3665 : Intrinsic::x86_avx512_mask_expand;
3666 Rep = Builder.CreateIntrinsic(
3668 }
else if (Name.starts_with(
"xop.vpcom")) {
3670 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3671 Name.ends_with(
"uq"))
3673 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3674 Name.ends_with(
"d") || Name.ends_with(
"q"))
3683 Name = Name.substr(9);
3684 if (Name.starts_with(
"lt"))
3686 else if (Name.starts_with(
"le"))
3688 else if (Name.starts_with(
"gt"))
3690 else if (Name.starts_with(
"ge"))
3692 else if (Name.starts_with(
"eq"))
3694 else if (Name.starts_with(
"ne"))
3696 else if (Name.starts_with(
"false"))
3698 else if (Name.starts_with(
"true"))
3705 }
else if (Name.starts_with(
"xop.vpcmov")) {
3707 Value *NotSel = Builder.CreateNot(Sel);
3710 Rep = Builder.CreateOr(Sel0, Sel1);
3711 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3712 Name.starts_with(
"avx512.mask.prol")) {
3714 }
else if (Name.starts_with(
"avx512.pror") ||
3715 Name.starts_with(
"avx512.mask.pror")) {
3717 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3718 Name.starts_with(
"avx512.mask.vpshld") ||
3719 Name.starts_with(
"avx512.maskz.vpshld")) {
3720 bool ZeroMask = Name[11] ==
'z';
3722 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3723 Name.starts_with(
"avx512.mask.vpshrd") ||
3724 Name.starts_with(
"avx512.maskz.vpshrd")) {
3725 bool ZeroMask = Name[11] ==
'z';
3727 }
else if (Name ==
"sse42.crc32.64.8") {
3730 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3732 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3733 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3734 Name.starts_with(
"avx512.vbroadcast.s")) {
3737 Type *EltTy = VecTy->getElementType();
3738 unsigned EltNum = VecTy->getNumElements();
3742 for (
unsigned I = 0;
I < EltNum; ++
I)
3743 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3744 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3745 Name.starts_with(
"sse41.pmovzx") ||
3746 Name.starts_with(
"avx2.pmovsx") ||
3747 Name.starts_with(
"avx2.pmovzx") ||
3748 Name.starts_with(
"avx512.mask.pmovsx") ||
3749 Name.starts_with(
"avx512.mask.pmovzx")) {
3751 unsigned NumDstElts = DstTy->getNumElements();
3755 for (
unsigned i = 0; i != NumDstElts; ++i)
3760 bool DoSext = Name.contains(
"pmovsx");
3762 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3767 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3768 Name ==
"avx512.mask.pmov.qd.512" ||
3769 Name ==
"avx512.mask.pmov.wb.256" ||
3770 Name ==
"avx512.mask.pmov.wb.512") {
3775 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3776 Name ==
"avx2.vbroadcasti128") {
3782 if (NumSrcElts == 2)
3785 Rep = Builder.CreateShuffleVector(
Load,
3787 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3788 Name.starts_with(
"avx512.mask.shuf.f")) {
3793 unsigned ControlBitsMask = NumLanes - 1;
3794 unsigned NumControlBits = NumLanes / 2;
3797 for (
unsigned l = 0; l != NumLanes; ++l) {
3798 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3800 if (l >= NumLanes / 2)
3801 LaneMask += NumLanes;
3802 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3803 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3809 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3810 Name.starts_with(
"avx512.mask.broadcasti")) {
3813 unsigned NumDstElts =
3817 for (
unsigned i = 0; i != NumDstElts; ++i)
3818 ShuffleMask[i] = i % NumSrcElts;
3824 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3825 Name.starts_with(
"avx2.vbroadcast") ||
3826 Name.starts_with(
"avx512.pbroadcast") ||
3827 Name.starts_with(
"avx512.mask.broadcast.s")) {
3834 Rep = Builder.CreateShuffleVector(
Op, M);
3839 }
else if (Name.starts_with(
"sse2.padds.") ||
3840 Name.starts_with(
"avx2.padds.") ||
3841 Name.starts_with(
"avx512.padds.") ||
3842 Name.starts_with(
"avx512.mask.padds.")) {
3844 }
else if (Name.starts_with(
"sse2.psubs.") ||
3845 Name.starts_with(
"avx2.psubs.") ||
3846 Name.starts_with(
"avx512.psubs.") ||
3847 Name.starts_with(
"avx512.mask.psubs.")) {
3849 }
else if (Name.starts_with(
"sse2.paddus.") ||
3850 Name.starts_with(
"avx2.paddus.") ||
3851 Name.starts_with(
"avx512.mask.paddus.")) {
3853 }
else if (Name.starts_with(
"sse2.psubus.") ||
3854 Name.starts_with(
"avx2.psubus.") ||
3855 Name.starts_with(
"avx512.mask.psubus.")) {
3857 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3862 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3866 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3871 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3876 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3877 Name ==
"avx512.psll.dq.512") {
3881 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3882 Name ==
"avx512.psrl.dq.512") {
3886 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3887 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3888 Name.starts_with(
"avx2.pblendd.")) {
3893 unsigned NumElts = VecTy->getNumElements();
3896 for (
unsigned i = 0; i != NumElts; ++i)
3897 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
3899 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3900 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3901 Name ==
"avx2.vinserti128" ||
3902 Name.starts_with(
"avx512.mask.insert")) {
3906 unsigned DstNumElts =
3908 unsigned SrcNumElts =
3910 unsigned Scale = DstNumElts / SrcNumElts;
3917 for (
unsigned i = 0; i != SrcNumElts; ++i)
3919 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3920 Idxs[i] = SrcNumElts;
3921 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3935 for (
unsigned i = 0; i != DstNumElts; ++i)
3938 for (
unsigned i = 0; i != SrcNumElts; ++i)
3939 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
3940 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3946 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3947 Name ==
"avx2.vextracti128" ||
3948 Name.starts_with(
"avx512.mask.vextract")) {
3951 unsigned DstNumElts =
3953 unsigned SrcNumElts =
3955 unsigned Scale = SrcNumElts / DstNumElts;
3962 for (
unsigned i = 0; i != DstNumElts; ++i) {
3963 Idxs[i] = i + (
Imm * DstNumElts);
3965 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3971 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3972 Name.starts_with(
"avx512.mask.perm.di.")) {
3976 unsigned NumElts = VecTy->getNumElements();
3979 for (
unsigned i = 0; i != NumElts; ++i)
3980 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
3982 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3987 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
3999 unsigned HalfSize = NumElts / 2;
4011 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4012 for (
unsigned i = 0; i < HalfSize; ++i)
4013 ShuffleMask[i] = StartIndex + i;
4016 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4017 for (
unsigned i = 0; i < HalfSize; ++i)
4018 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4020 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4022 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4023 Name.starts_with(
"avx512.mask.vpermil.p") ||
4024 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4028 unsigned NumElts = VecTy->getNumElements();
4030 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4031 unsigned IdxMask = ((1 << IdxSize) - 1);
4037 for (
unsigned i = 0; i != NumElts; ++i)
4038 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4040 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4045 }
else if (Name ==
"sse2.pshufl.w" ||
4046 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4051 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4055 for (
unsigned l = 0; l != NumElts; l += 8) {
4056 for (
unsigned i = 0; i != 4; ++i)
4057 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4058 for (
unsigned i = 4; i != 8; ++i)
4059 Idxs[i + l] = i + l;
4062 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4067 }
else if (Name ==
"sse2.pshufh.w" ||
4068 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4073 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4077 for (
unsigned l = 0; l != NumElts; l += 8) {
4078 for (
unsigned i = 0; i != 4; ++i)
4079 Idxs[i + l] = i + l;
4080 for (
unsigned i = 0; i != 4; ++i)
4081 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4084 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4089 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4096 unsigned HalfLaneElts = NumLaneElts / 2;
4099 for (
unsigned i = 0; i != NumElts; ++i) {
4101 Idxs[i] = i - (i % NumLaneElts);
4103 if ((i % NumLaneElts) >= HalfLaneElts)
4107 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4110 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4114 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4115 Name.starts_with(
"avx512.mask.movshdup") ||
4116 Name.starts_with(
"avx512.mask.movsldup")) {
4122 if (Name.starts_with(
"avx512.mask.movshdup."))
4126 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4127 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4128 Idxs[i + l + 0] = i + l +
Offset;
4129 Idxs[i + l + 1] = i + l +
Offset;
4132 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4136 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4137 Name.starts_with(
"avx512.mask.unpckl.")) {
4144 for (
int l = 0; l != NumElts; l += NumLaneElts)
4145 for (
int i = 0; i != NumLaneElts; ++i)
4146 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4148 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4152 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4153 Name.starts_with(
"avx512.mask.unpckh.")) {
4160 for (
int l = 0; l != NumElts; l += NumLaneElts)
4161 for (
int i = 0; i != NumLaneElts; ++i)
4162 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4164 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4168 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4169 Name.starts_with(
"avx512.mask.pand.")) {
4172 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4174 Rep = Builder.CreateBitCast(Rep, FTy);
4177 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4178 Name.starts_with(
"avx512.mask.pandn.")) {
4181 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4182 Rep = Builder.CreateAnd(Rep,
4184 Rep = Builder.CreateBitCast(Rep, FTy);
4187 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4188 Name.starts_with(
"avx512.mask.por.")) {
4191 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4193 Rep = Builder.CreateBitCast(Rep, FTy);
4196 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4197 Name.starts_with(
"avx512.mask.pxor.")) {
4200 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4202 Rep = Builder.CreateBitCast(Rep, FTy);
4205 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4209 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4213 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4217 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4218 if (Name.ends_with(
".512")) {
4220 if (Name[17] ==
's')
4221 IID = Intrinsic::x86_avx512_add_ps_512;
4223 IID = Intrinsic::x86_avx512_add_pd_512;
4225 Rep = Builder.CreateIntrinsic(
4233 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4234 if (Name.ends_with(
".512")) {
4236 if (Name[17] ==
's')
4237 IID = Intrinsic::x86_avx512_div_ps_512;
4239 IID = Intrinsic::x86_avx512_div_pd_512;
4241 Rep = Builder.CreateIntrinsic(
4249 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4250 if (Name.ends_with(
".512")) {
4252 if (Name[17] ==
's')
4253 IID = Intrinsic::x86_avx512_mul_ps_512;
4255 IID = Intrinsic::x86_avx512_mul_pd_512;
4257 Rep = Builder.CreateIntrinsic(
4265 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4266 if (Name.ends_with(
".512")) {
4268 if (Name[17] ==
's')
4269 IID = Intrinsic::x86_avx512_sub_ps_512;
4271 IID = Intrinsic::x86_avx512_sub_pd_512;
4273 Rep = Builder.CreateIntrinsic(
4281 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4282 Name.starts_with(
"avx512.mask.min.p")) &&
4283 Name.drop_front(18) ==
".512") {
4284 bool IsDouble = Name[17] ==
'd';
4285 bool IsMin = Name[13] ==
'i';
4287 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4288 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4291 Rep = Builder.CreateIntrinsic(
4296 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4298 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4299 {CI->getArgOperand(0), Builder.getInt1(false)});
4302 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4303 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4304 bool IsVariable = Name[16] ==
'v';
4305 char Size = Name[16] ==
'.' ? Name[17]
4306 : Name[17] ==
'.' ? Name[18]
4307 : Name[18] ==
'.' ? Name[19]
4311 if (IsVariable && Name[17] !=
'.') {
4312 if (
Size ==
'd' && Name[17] ==
'2')
4313 IID = Intrinsic::x86_avx2_psllv_q;
4314 else if (
Size ==
'd' && Name[17] ==
'4')
4315 IID = Intrinsic::x86_avx2_psllv_q_256;
4316 else if (
Size ==
's' && Name[17] ==
'4')
4317 IID = Intrinsic::x86_avx2_psllv_d;
4318 else if (
Size ==
's' && Name[17] ==
'8')
4319 IID = Intrinsic::x86_avx2_psllv_d_256;
4320 else if (
Size ==
'h' && Name[17] ==
'8')
4321 IID = Intrinsic::x86_avx512_psllv_w_128;
4322 else if (
Size ==
'h' && Name[17] ==
'1')
4323 IID = Intrinsic::x86_avx512_psllv_w_256;
4324 else if (Name[17] ==
'3' && Name[18] ==
'2')
4325 IID = Intrinsic::x86_avx512_psllv_w_512;
4328 }
else if (Name.ends_with(
".128")) {
4330 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4331 : Intrinsic::x86_sse2_psll_d;
4332 else if (
Size ==
'q')
4333 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4334 : Intrinsic::x86_sse2_psll_q;
4335 else if (
Size ==
'w')
4336 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4337 : Intrinsic::x86_sse2_psll_w;
4340 }
else if (Name.ends_with(
".256")) {
4342 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4343 : Intrinsic::x86_avx2_psll_d;
4344 else if (
Size ==
'q')
4345 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4346 : Intrinsic::x86_avx2_psll_q;
4347 else if (
Size ==
'w')
4348 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4349 : Intrinsic::x86_avx2_psll_w;
4354 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4355 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4356 : Intrinsic::x86_avx512_psll_d_512;
4357 else if (
Size ==
'q')
4358 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4359 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4360 : Intrinsic::x86_avx512_psll_q_512;
4361 else if (
Size ==
'w')
4362 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4363 : Intrinsic::x86_avx512_psll_w_512;
4369 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4370 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4371 bool IsVariable = Name[16] ==
'v';
4372 char Size = Name[16] ==
'.' ? Name[17]
4373 : Name[17] ==
'.' ? Name[18]
4374 : Name[18] ==
'.' ? Name[19]
4378 if (IsVariable && Name[17] !=
'.') {
4379 if (
Size ==
'd' && Name[17] ==
'2')
4380 IID = Intrinsic::x86_avx2_psrlv_q;
4381 else if (
Size ==
'd' && Name[17] ==
'4')
4382 IID = Intrinsic::x86_avx2_psrlv_q_256;
4383 else if (
Size ==
's' && Name[17] ==
'4')
4384 IID = Intrinsic::x86_avx2_psrlv_d;
4385 else if (
Size ==
's' && Name[17] ==
'8')
4386 IID = Intrinsic::x86_avx2_psrlv_d_256;
4387 else if (
Size ==
'h' && Name[17] ==
'8')
4388 IID = Intrinsic::x86_avx512_psrlv_w_128;
4389 else if (
Size ==
'h' && Name[17] ==
'1')
4390 IID = Intrinsic::x86_avx512_psrlv_w_256;
4391 else if (Name[17] ==
'3' && Name[18] ==
'2')
4392 IID = Intrinsic::x86_avx512_psrlv_w_512;
4395 }
else if (Name.ends_with(
".128")) {
4397 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4398 : Intrinsic::x86_sse2_psrl_d;
4399 else if (
Size ==
'q')
4400 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4401 : Intrinsic::x86_sse2_psrl_q;
4402 else if (
Size ==
'w')
4403 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4404 : Intrinsic::x86_sse2_psrl_w;
4407 }
else if (Name.ends_with(
".256")) {
4409 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4410 : Intrinsic::x86_avx2_psrl_d;
4411 else if (
Size ==
'q')
4412 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4413 : Intrinsic::x86_avx2_psrl_q;
4414 else if (
Size ==
'w')
4415 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4416 : Intrinsic::x86_avx2_psrl_w;
4421 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4422 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4423 : Intrinsic::x86_avx512_psrl_d_512;
4424 else if (
Size ==
'q')
4425 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4426 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4427 : Intrinsic::x86_avx512_psrl_q_512;
4428 else if (
Size ==
'w')
4429 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4430 : Intrinsic::x86_avx512_psrl_w_512;
4436 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4437 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4438 bool IsVariable = Name[16] ==
'v';
4439 char Size = Name[16] ==
'.' ? Name[17]
4440 : Name[17] ==
'.' ? Name[18]
4441 : Name[18] ==
'.' ? Name[19]
4445 if (IsVariable && Name[17] !=
'.') {
4446 if (
Size ==
's' && Name[17] ==
'4')
4447 IID = Intrinsic::x86_avx2_psrav_d;
4448 else if (
Size ==
's' && Name[17] ==
'8')
4449 IID = Intrinsic::x86_avx2_psrav_d_256;
4450 else if (
Size ==
'h' && Name[17] ==
'8')
4451 IID = Intrinsic::x86_avx512_psrav_w_128;
4452 else if (
Size ==
'h' && Name[17] ==
'1')
4453 IID = Intrinsic::x86_avx512_psrav_w_256;
4454 else if (Name[17] ==
'3' && Name[18] ==
'2')
4455 IID = Intrinsic::x86_avx512_psrav_w_512;
4458 }
else if (Name.ends_with(
".128")) {
4460 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4461 : Intrinsic::x86_sse2_psra_d;
4462 else if (
Size ==
'q')
4463 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4464 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4465 : Intrinsic::x86_avx512_psra_q_128;
4466 else if (
Size ==
'w')
4467 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4468 : Intrinsic::x86_sse2_psra_w;
4471 }
else if (Name.ends_with(
".256")) {
4473 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4474 : Intrinsic::x86_avx2_psra_d;
4475 else if (
Size ==
'q')
4476 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4477 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4478 : Intrinsic::x86_avx512_psra_q_256;
4479 else if (
Size ==
'w')
4480 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4481 : Intrinsic::x86_avx2_psra_w;
4486 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4487 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4488 : Intrinsic::x86_avx512_psra_d_512;
4489 else if (
Size ==
'q')
4490 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4491 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4492 : Intrinsic::x86_avx512_psra_q_512;
4493 else if (
Size ==
'w')
4494 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4495 : Intrinsic::x86_avx512_psra_w_512;
4501 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4503 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4505 }
else if (Name.ends_with(
".movntdqa")) {
4509 LoadInst *LI = Builder.CreateAlignedLoad(
4514 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4515 Name.starts_with(
"fma.vfmsub.") ||
4516 Name.starts_with(
"fma.vfnmadd.") ||
4517 Name.starts_with(
"fma.vfnmsub.")) {
4518 bool NegMul = Name[6] ==
'n';
4519 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4520 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4531 if (NegMul && !IsScalar)
4532 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4533 if (NegMul && IsScalar)
4534 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4536 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4538 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4542 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4550 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4554 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4555 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4556 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4557 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4558 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4559 bool IsMask3 = Name[11] ==
'3';
4560 bool IsMaskZ = Name[11] ==
'z';
4562 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4563 bool NegMul = Name[2] ==
'n';
4564 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4570 if (NegMul && (IsMask3 || IsMaskZ))
4571 A = Builder.CreateFNeg(
A);
4572 if (NegMul && !(IsMask3 || IsMaskZ))
4573 B = Builder.CreateFNeg(
B);
4575 C = Builder.CreateFNeg(
C);
4577 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4578 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4579 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4586 if (Name.back() ==
'd')
4587 IID = Intrinsic::x86_avx512_vfmadd_f64;
4589 IID = Intrinsic::x86_avx512_vfmadd_f32;
4590 Rep = Builder.CreateIntrinsic(IID,
Ops);
4592 Rep = Builder.CreateFMA(
A,
B,
C);
4601 if (NegAcc && IsMask3)
4606 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4608 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4609 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4610 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4611 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4612 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4613 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4614 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4615 bool IsMask3 = Name[11] ==
'3';
4616 bool IsMaskZ = Name[11] ==
'z';
4618 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4619 bool NegMul = Name[2] ==
'n';
4620 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4626 if (NegMul && (IsMask3 || IsMaskZ))
4627 A = Builder.CreateFNeg(
A);
4628 if (NegMul && !(IsMask3 || IsMaskZ))
4629 B = Builder.CreateFNeg(
B);
4631 C = Builder.CreateFNeg(
C);
4638 if (Name[Name.size() - 5] ==
's')
4639 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4641 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4645 Rep = Builder.CreateFMA(
A,
B,
C);
4653 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4657 if (VecWidth == 128 && EltWidth == 32)
4658 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4659 else if (VecWidth == 256 && EltWidth == 32)
4660 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4661 else if (VecWidth == 128 && EltWidth == 64)
4662 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4663 else if (VecWidth == 256 && EltWidth == 64)
4664 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4670 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4671 Rep = Builder.CreateIntrinsic(IID,
Ops);
4672 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4673 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4674 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4675 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4676 bool IsMask3 = Name[11] ==
'3';
4677 bool IsMaskZ = Name[11] ==
'z';
4679 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4680 bool IsSubAdd = Name[3] ==
's';
4684 if (Name[Name.size() - 5] ==
's')
4685 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4687 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4692 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4694 Rep = Builder.CreateIntrinsic(IID,
Ops);
4703 Value *Odd = Builder.CreateCall(FMA,
Ops);
4704 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4705 Value *Even = Builder.CreateCall(FMA,
Ops);
4711 for (
int i = 0; i != NumElts; ++i)
4712 Idxs[i] = i + (i % 2) * NumElts;
4714 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4722 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4723 Name.starts_with(
"avx512.maskz.pternlog.")) {
4724 bool ZeroMask = Name[11] ==
'z';
4728 if (VecWidth == 128 && EltWidth == 32)
4729 IID = Intrinsic::x86_avx512_pternlog_d_128;
4730 else if (VecWidth == 256 && EltWidth == 32)
4731 IID = Intrinsic::x86_avx512_pternlog_d_256;
4732 else if (VecWidth == 512 && EltWidth == 32)
4733 IID = Intrinsic::x86_avx512_pternlog_d_512;
4734 else if (VecWidth == 128 && EltWidth == 64)
4735 IID = Intrinsic::x86_avx512_pternlog_q_128;
4736 else if (VecWidth == 256 && EltWidth == 64)
4737 IID = Intrinsic::x86_avx512_pternlog_q_256;
4738 else if (VecWidth == 512 && EltWidth == 64)
4739 IID = Intrinsic::x86_avx512_pternlog_q_512;
4745 Rep = Builder.CreateIntrinsic(IID, Args);
4749 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4750 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4751 bool ZeroMask = Name[11] ==
'z';
4752 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4755 if (VecWidth == 128 && !
High)
4756 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4757 else if (VecWidth == 256 && !
High)
4758 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4759 else if (VecWidth == 512 && !
High)
4760 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4761 else if (VecWidth == 128 &&
High)
4762 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4763 else if (VecWidth == 256 &&
High)
4764 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4765 else if (VecWidth == 512 &&
High)
4766 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4772 Rep = Builder.CreateIntrinsic(IID, Args);
4776 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4777 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4778 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4779 bool ZeroMask = Name[11] ==
'z';
4780 bool IndexForm = Name[17] ==
'i';
4782 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4783 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4784 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4785 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4786 bool ZeroMask = Name[11] ==
'z';
4787 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4790 if (VecWidth == 128 && !IsSaturating)
4791 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4792 else if (VecWidth == 256 && !IsSaturating)
4793 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4794 else if (VecWidth == 512 && !IsSaturating)
4795 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4796 else if (VecWidth == 128 && IsSaturating)
4797 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4798 else if (VecWidth == 256 && IsSaturating)
4799 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4800 else if (VecWidth == 512 && IsSaturating)
4801 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4811 if (Args[1]->
getType()->isVectorTy() &&
4814 ->isIntegerTy(32) &&
4815 Args[2]->
getType()->isVectorTy() &&
4818 ->isIntegerTy(32)) {
4819 Type *NewArgType =
nullptr;
4820 if (VecWidth == 128)
4822 else if (VecWidth == 256)
4824 else if (VecWidth == 512)
4830 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4831 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4834 Rep = Builder.CreateIntrinsic(IID, Args);
4838 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4839 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4840 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4841 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4842 bool ZeroMask = Name[11] ==
'z';
4843 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4846 if (VecWidth == 128 && !IsSaturating)
4847 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4848 else if (VecWidth == 256 && !IsSaturating)
4849 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4850 else if (VecWidth == 512 && !IsSaturating)
4851 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4852 else if (VecWidth == 128 && IsSaturating)
4853 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4854 else if (VecWidth == 256 && IsSaturating)
4855 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4856 else if (VecWidth == 512 && IsSaturating)
4857 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4867 if (Args[1]->
getType()->isVectorTy() &&
4870 ->isIntegerTy(32) &&
4871 Args[2]->
getType()->isVectorTy() &&
4874 ->isIntegerTy(32)) {
4875 Type *NewArgType =
nullptr;
4876 if (VecWidth == 128)
4878 else if (VecWidth == 256)
4880 else if (VecWidth == 512)
4886 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4887 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4890 Rep = Builder.CreateIntrinsic(IID, Args);
4894 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4895 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4896 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4898 if (Name[0] ==
'a' && Name.back() ==
'2')
4899 IID = Intrinsic::x86_addcarry_32;
4900 else if (Name[0] ==
'a' && Name.back() ==
'4')
4901 IID = Intrinsic::x86_addcarry_64;
4902 else if (Name[0] ==
's' && Name.back() ==
'2')
4903 IID = Intrinsic::x86_subborrow_32;
4904 else if (Name[0] ==
's' && Name.back() ==
'4')
4905 IID = Intrinsic::x86_subborrow_64;
4912 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4915 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4918 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4922 }
else if (Name.starts_with(
"avx512.mask.") &&
4925 }
else if (Name.starts_with(
"bmi.pdep.")) {
4927 }
else if (Name.starts_with(
"bmi.pext.")) {
4937 if (Name.starts_with(
"neon.bfcvt")) {
4938 if (Name.starts_with(
"neon.bfcvtn2")) {
4940 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4942 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4943 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4946 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4947 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4949 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4953 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4954 return Builder.CreateShuffleVector(
4957 return Builder.CreateFPTrunc(CI->
getOperand(0),
4960 }
else if (Name.starts_with(
"sve.fcvt")) {
4963 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4964 .
Case(
"sve.fcvtnt.bf16f32",
4965 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4977 if (Args[1]->
getType() != BadPredTy)
4980 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
4981 BadPredTy, Args[1]);
4982 Args[1] = Builder.CreateIntrinsic(
4983 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
4985 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
4989 if (Name ==
"neon.vcvtfp2hf")
4990 return Builder.CreateBitCast(
4991 Builder.CreateFPTrunc(
4995 if (Name ==
"neon.vcvthf2fp")
4996 return Builder.CreateFPExt(
4997 Builder.CreateBitCast(
5007 if (Name ==
"mve.vctp64.old") {
5010 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5013 Value *C1 = Builder.CreateIntrinsic(
5014 Intrinsic::arm_mve_pred_v2i,
5016 return Builder.CreateIntrinsic(
5017 Intrinsic::arm_mve_pred_i2v,
5019 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5020 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5021 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5022 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5024 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5025 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5026 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5027 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5029 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5030 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5031 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5032 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5033 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5034 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5035 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5036 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5037 std::vector<Type *> Tys;
5041 case Intrinsic::arm_mve_mull_int_predicated:
5042 case Intrinsic::arm_mve_vqdmull_predicated:
5043 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5046 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5047 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5048 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5052 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5056 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5060 case Intrinsic::arm_cde_vcx1q_predicated:
5061 case Intrinsic::arm_cde_vcx1qa_predicated:
5062 case Intrinsic::arm_cde_vcx2q_predicated:
5063 case Intrinsic::arm_cde_vcx2qa_predicated:
5064 case Intrinsic::arm_cde_vcx3q_predicated:
5065 case Intrinsic::arm_cde_vcx3qa_predicated:
5072 std::vector<Value *>
Ops;
5074 Type *Ty =
Op->getType();
5075 if (Ty->getScalarSizeInBits() == 1) {
5076 Value *C1 = Builder.CreateIntrinsic(
5077 Intrinsic::arm_mve_pred_v2i,
5079 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5084 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5099 auto UpgradeLegacyWMMAIUIntrinsicCall =
5104 Args.push_back(Builder.getFalse());
5108 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5115 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5119 NewCall->copyMetadata(*CI);
5123 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5124 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5125 "intrinsic should have 7 arguments");
5128 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5130 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5131 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5132 "intrinsic should have 8 arguments");
5137 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5140 switch (
F->getIntrinsicID()) {
5143 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5144 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5145 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5146 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5147 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5148 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5163 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5166 F->getParent(),
F->getIntrinsicID(), Overloads);
5171 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5175 NewCall->copyMetadata(*CI);
5176 NewCall->takeName(CI);
5181 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5187 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5188 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5213 if (NumOperands < 3)
5226 bool IsVolatile =
false;
5230 if (NumOperands > 3)
5235 if (NumOperands > 5) {
5237 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5251 if (VT->getElementType()->isIntegerTy(16)) {
5254 Val = Builder.CreateBitCast(Val, AsBF16);
5262 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5264 unsigned AddrSpace = PtrTy->getAddressSpace();
5267 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5269 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5274 MDNode *RangeNotPrivate =
5277 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5283 return Builder.CreateBitCast(RMW, RetTy);
5304 return MAV->getMetadata();
5313 if (Name ==
"label") {
5315 }
else if (Name ==
"assign") {
5322 }
else if (Name ==
"declare") {
5326 }
else if (Name ==
"addr") {
5336 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5337 }
else if (Name ==
"value") {
5340 unsigned ExprOp = 2;
5355 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5363 int64_t OffsetVal =
Offset->getSExtValue();
5364 return Builder.CreateIntrinsic(OffsetVal >= 0
5365 ? Intrinsic::vector_splice_left
5366 : Intrinsic::vector_splice_right,
5368 {CI->getArgOperand(0), CI->getArgOperand(1),
5369 Builder.getInt32(std::abs(OffsetVal))});
5374 if (Name.starts_with(
"to.fp16")) {
5376 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5377 return Builder.CreateBitCast(Cast, CI->
getType());
5380 if (Name.starts_with(
"from.fp16")) {
5382 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5383 return Builder.CreateFPExt(Cast, CI->
getType());
5442 else if (Opcode == Instruction::ICmp)
5445 else if (Opcode == Instruction::FCmp)
5448 else if (Opcode == Instruction::Select)
5453 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5465 if (Defaults.empty())
5468 unsigned OldArgCount = CI->
arg_size();
5469 unsigned NewArgCount = NewFn->
arg_size();
5473 if (OldArgCount >= NewArgCount)
5481 if (OldArgCount < FirstDefault)
5486 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5487 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5488 "missing argument outside the default range");
5489 Type *ParamTy = NewFT->getParamType(Idx);
5494 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5500 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5532 if (!Name.consume_front(
"llvm."))
5535 bool IsX86 = Name.consume_front(
"x86.");
5536 bool IsNVVM = Name.consume_front(
"nvvm.");
5537 bool IsAArch64 = Name.consume_front(
"aarch64.");
5538 bool IsARM = Name.consume_front(
"arm.");
5539 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5540 bool IsDbg = Name.consume_front(
"dbg.");
5542 (Name.consume_front(
"experimental.vector.splice") ||
5543 Name.consume_front(
"vector.splice")) &&
5544 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5545 Value *Rep =
nullptr;
5547 if (!IsX86 && Name ==
"stackprotectorcheck") {
5549 }
else if (IsNVVM) {
5553 }
else if (IsAArch64) {
5557 }
else if (IsAMDGCN) {
5561 }
else if (IsOldSplice) {
5563 }
else if (Name.consume_front(
"convert.")) {
5565 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5580 const auto &DefaultCase = [&]() ->
void {
5588 "Unknown function for CallBase upgrade and isn't just a name change");
5596 "Return type must have changed");
5597 assert(OldST->getNumElements() ==
5599 "Must have same number of elements");
5602 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5605 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5606 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5607 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5631 case Intrinsic::arm_neon_vst1:
5632 case Intrinsic::arm_neon_vst2:
5633 case Intrinsic::arm_neon_vst3:
5634 case Intrinsic::arm_neon_vst4:
5635 case Intrinsic::arm_neon_vst2lane:
5636 case Intrinsic::arm_neon_vst3lane:
5637 case Intrinsic::arm_neon_vst4lane: {
5639 NewCall = Builder.CreateCall(NewFn, Args);
5642 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5643 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5644 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5649 NewCall = Builder.CreateCall(NewFn, Args);
5652 case Intrinsic::aarch64_sve_ld3_sret:
5653 case Intrinsic::aarch64_sve_ld4_sret:
5654 case Intrinsic::aarch64_sve_ld2_sret: {
5662 Name = Name.substr(5);
5669 unsigned MinElts = RetTy->getMinNumElements() /
N;
5671 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5673 for (
unsigned I = 0;
I <
N;
I++) {
5674 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5675 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5681 case Intrinsic::coro_end_async:
5682 case Intrinsic::coro_end: {
5684 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5686 NewCall = Builder.CreateCall(NewFn, Args);
5691 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5692 Value *InRamp = Builder.CreateCall(IsInRamp);
5702 case Intrinsic::vector_extract: {
5704 Name = Name.substr(5);
5705 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5710 unsigned MinElts = RetTy->getMinNumElements();
5713 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5717 case Intrinsic::vector_insert: {
5719 Name = Name.substr(5);
5720 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5724 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5729 NewCall = Builder.CreateCall(
5733 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5739 assert(
N > 1 &&
"Create is expected to be between 2-4");
5742 unsigned MinElts = RetTy->getMinNumElements() /
N;
5743 for (
unsigned I = 0;
I <
N;
I++) {
5745 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5752 case Intrinsic::arm_neon_bfdot:
5753 case Intrinsic::arm_neon_bfmmla:
5754 case Intrinsic::arm_neon_bfmlalb:
5755 case Intrinsic::arm_neon_bfmlalt:
5756 case Intrinsic::aarch64_neon_bfdot:
5757 case Intrinsic::aarch64_neon_bfmmla:
5758 case Intrinsic::aarch64_neon_bfmlalb:
5759 case Intrinsic::aarch64_neon_bfmlalt: {
5762 "Mismatch between function args and call args");
5763 size_t OperandWidth =
5765 assert((OperandWidth == 64 || OperandWidth == 128) &&
5766 "Unexpected operand width");
5768 auto Iter = CI->
args().begin();
5769 Args.push_back(*Iter++);
5770 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5771 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5772 NewCall = Builder.CreateCall(NewFn, Args);
5776 case Intrinsic::bitreverse:
5777 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5780 case Intrinsic::ctlz:
5781 case Intrinsic::cttz: {
5788 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5792 case Intrinsic::objectsize: {
5793 Value *NullIsUnknownSize =
5797 NewCall = Builder.CreateCall(
5802 case Intrinsic::ctpop:
5803 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5805 case Intrinsic::dbg_value: {
5807 Name = Name.substr(5);
5809 if (Name.starts_with(
"dbg.addr")) {
5823 if (
Offset->isNullValue()) {
5824 NewCall = Builder.CreateCall(
5833 case Intrinsic::ptr_annotation:
5841 NewCall = Builder.CreateCall(
5850 case Intrinsic::var_annotation:
5857 NewCall = Builder.CreateCall(
5866 case Intrinsic::riscv_aes32dsi:
5867 case Intrinsic::riscv_aes32dsmi:
5868 case Intrinsic::riscv_aes32esi:
5869 case Intrinsic::riscv_aes32esmi:
5870 case Intrinsic::riscv_sm4ks:
5871 case Intrinsic::riscv_sm4ed: {
5881 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5882 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5888 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5889 Value *Res = NewCall;
5891 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5897 case Intrinsic::nvvm_mapa_shared_cluster: {
5901 Value *Res = NewCall;
5902 Res = Builder.CreateAddrSpaceCast(
5909 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5910 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5913 Args[0] = Builder.CreateAddrSpaceCast(
5916 NewCall = Builder.CreateCall(NewFn, Args);
5922 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5923 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5924 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5925 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5926 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5927 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5928 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5929 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5936 Args[0] = Builder.CreateAddrSpaceCast(
5945 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5947 NewCall = Builder.CreateCall(NewFn, Args);
5953 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
5954 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
5955 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
5956 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
5957 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
5958 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
5959 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
5960 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
5962 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
5966 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
5967 NewCall = Builder.CreateCall(NewFn, Args);
5970 case Intrinsic::nvvm_tcgen05_mma_shared:
5971 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5972 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5973 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
5974 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
5975 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
5976 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
5977 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
5978 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
5979 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
5980 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
5981 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
5982 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
5983 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
5984 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
5985 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
5986 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
5987 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
5988 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
5989 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
5990 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
5991 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
5992 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
5993 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
5994 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
5995 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
5996 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
5997 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
5998 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
5999 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
6000 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
6001 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
6002 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
6003 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
6004 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
6005 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
6006 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
6007 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
6008 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
6010 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
6011 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
6013 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
6014 case Intrinsic::nvvm_tcgen05_mma_tensor:
6015 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
6016 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
6017 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
6018 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
6019 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
6020 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
6021 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
6022 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
6023 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
6024 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
6025 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
6026 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
6027 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
6028 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
6030 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
6031 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
6033 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
6035 Args.push_back(Builder.getInt32(0));
6036 NewCall = Builder.CreateCall(NewFn, Args);
6039 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6040 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6041 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6042 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6045 Builder.getFalse()});
6047 case Intrinsic::riscv_sha256sig0:
6048 case Intrinsic::riscv_sha256sig1:
6049 case Intrinsic::riscv_sha256sum0:
6050 case Intrinsic::riscv_sha256sum1:
6051 case Intrinsic::riscv_sm3p0:
6052 case Intrinsic::riscv_sm3p1: {
6059 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6061 NewCall = Builder.CreateCall(NewFn, Arg);
6063 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6070 case Intrinsic::x86_xop_vfrcz_ss:
6071 case Intrinsic::x86_xop_vfrcz_sd:
6072 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6075 case Intrinsic::x86_xop_vpermil2pd:
6076 case Intrinsic::x86_xop_vpermil2ps:
6077 case Intrinsic::x86_xop_vpermil2pd_256:
6078 case Intrinsic::x86_xop_vpermil2ps_256: {
6082 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6083 NewCall = Builder.CreateCall(NewFn, Args);
6087 case Intrinsic::x86_sse41_ptestc:
6088 case Intrinsic::x86_sse41_ptestz:
6089 case Intrinsic::x86_sse41_ptestnzc: {
6103 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6104 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6106 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6110 case Intrinsic::x86_rdtscp: {
6116 NewCall = Builder.CreateCall(NewFn);
6118 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6121 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6129 case Intrinsic::x86_sse41_insertps:
6130 case Intrinsic::x86_sse41_dppd:
6131 case Intrinsic::x86_sse41_dpps:
6132 case Intrinsic::x86_sse41_mpsadbw:
6133 case Intrinsic::x86_avx_dp_ps_256:
6134 case Intrinsic::x86_avx2_mpsadbw: {
6140 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6141 NewCall = Builder.CreateCall(NewFn, Args);
6145 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6146 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6147 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6148 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6149 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6150 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6156 NewCall = Builder.CreateCall(NewFn, Args);
6165 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6166 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6167 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6168 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6169 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6170 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6174 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6175 Args[1] = Builder.CreateBitCast(
6178 NewCall = Builder.CreateCall(NewFn, Args);
6179 Value *Res = Builder.CreateBitCast(
6187 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6188 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6189 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6193 Args[1] = Builder.CreateBitCast(
6195 Args[2] = Builder.CreateBitCast(
6198 NewCall = Builder.CreateCall(NewFn, Args);
6202 case Intrinsic::thread_pointer: {
6203 NewCall = Builder.CreateCall(NewFn, {});
6207 case Intrinsic::memcpy:
6208 case Intrinsic::memmove:
6209 case Intrinsic::memset: {
6225 NewCall = Builder.CreateCall(NewFn, Args);
6227 AttributeList NewAttrs = AttributeList::get(
6228 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6229 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6230 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6235 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6238 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6242 case Intrinsic::masked_load:
6243 case Intrinsic::masked_gather:
6244 case Intrinsic::masked_store:
6245 case Intrinsic::masked_scatter: {
6251 auto GetMaybeAlign = [](
Value *
Op) {
6253 uint64_t Val = CI->getZExtValue();
6261 auto GetAlign = [&](
Value *
Op) {
6270 case Intrinsic::masked_load:
6271 NewCall = Builder.CreateMaskedLoad(
6275 case Intrinsic::masked_gather:
6276 NewCall = Builder.CreateMaskedGather(
6282 case Intrinsic::masked_store:
6283 NewCall = Builder.CreateMaskedStore(
6287 case Intrinsic::masked_scatter:
6288 NewCall = Builder.CreateMaskedScatter(
6290 DL.getValueOrABITypeAlignment(
6304 case Intrinsic::lifetime_start:
6305 case Intrinsic::lifetime_end: {
6317 NewCall = Builder.CreateLifetimeStart(Ptr);
6319 NewCall = Builder.CreateLifetimeEnd(Ptr);
6328 case Intrinsic::x86_avx512_vpdpbusd_128:
6329 case Intrinsic::x86_avx512_vpdpbusd_256:
6330 case Intrinsic::x86_avx512_vpdpbusd_512:
6331 case Intrinsic::x86_avx512_vpdpbusds_128:
6332 case Intrinsic::x86_avx512_vpdpbusds_256:
6333 case Intrinsic::x86_avx512_vpdpbusds_512:
6334 case Intrinsic::x86_avx2_vpdpbssd_128:
6335 case Intrinsic::x86_avx2_vpdpbssd_256:
6336 case Intrinsic::x86_avx10_vpdpbssd_512:
6337 case Intrinsic::x86_avx2_vpdpbssds_128:
6338 case Intrinsic::x86_avx2_vpdpbssds_256:
6339 case Intrinsic::x86_avx10_vpdpbssds_512:
6340 case Intrinsic::x86_avx2_vpdpbsud_128:
6341 case Intrinsic::x86_avx2_vpdpbsud_256:
6342 case Intrinsic::x86_avx10_vpdpbsud_512:
6343 case Intrinsic::x86_avx2_vpdpbsuds_128:
6344 case Intrinsic::x86_avx2_vpdpbsuds_256:
6345 case Intrinsic::x86_avx10_vpdpbsuds_512:
6346 case Intrinsic::x86_avx2_vpdpbuud_128:
6347 case Intrinsic::x86_avx2_vpdpbuud_256:
6348 case Intrinsic::x86_avx10_vpdpbuud_512:
6349 case Intrinsic::x86_avx2_vpdpbuuds_128:
6350 case Intrinsic::x86_avx2_vpdpbuuds_256:
6351 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6356 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6357 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6359 NewCall = Builder.CreateCall(NewFn, Args);
6362 case Intrinsic::x86_avx512_vpdpwssd_128:
6363 case Intrinsic::x86_avx512_vpdpwssd_256:
6364 case Intrinsic::x86_avx512_vpdpwssd_512:
6365 case Intrinsic::x86_avx512_vpdpwssds_128:
6366 case Intrinsic::x86_avx512_vpdpwssds_256:
6367 case Intrinsic::x86_avx512_vpdpwssds_512:
6368 case Intrinsic::x86_avx2_vpdpwsud_128:
6369 case Intrinsic::x86_avx2_vpdpwsud_256:
6370 case Intrinsic::x86_avx10_vpdpwsud_512:
6371 case Intrinsic::x86_avx2_vpdpwsuds_128:
6372 case Intrinsic::x86_avx2_vpdpwsuds_256:
6373 case Intrinsic::x86_avx10_vpdpwsuds_512:
6374 case Intrinsic::x86_avx2_vpdpwusd_128:
6375 case Intrinsic::x86_avx2_vpdpwusd_256:
6376 case Intrinsic::x86_avx10_vpdpwusd_512:
6377 case Intrinsic::x86_avx2_vpdpwusds_128:
6378 case Intrinsic::x86_avx2_vpdpwusds_256:
6379 case Intrinsic::x86_avx10_vpdpwusds_512:
6380 case Intrinsic::x86_avx2_vpdpwuud_128:
6381 case Intrinsic::x86_avx2_vpdpwuud_256:
6382 case Intrinsic::x86_avx10_vpdpwuud_512:
6383 case Intrinsic::x86_avx2_vpdpwuuds_128:
6384 case Intrinsic::x86_avx2_vpdpwuuds_256:
6385 case Intrinsic::x86_avx10_vpdpwuuds_512:
6390 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6391 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6393 NewCall = Builder.CreateCall(NewFn, Args);
6396 assert(NewCall &&
"Should have either set this variable or returned through "
6397 "the default case");
6404 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6418 F->eraseFromParent();
6424 if (NumOperands == 0)
6432 if (NumOperands == 3) {
6436 Metadata *Elts2[] = {ScalarType, ScalarType,
6450 if (
Opc != Instruction::BitCast)
6454 Type *SrcTy = V->getType();
6471 if (
Opc != Instruction::BitCast)
6474 Type *SrcTy =
C->getType();
6491 if (Flag.getNumOperands() < 3)
6492 return std::nullopt;
6494 return Name->getString();
6495 return std::nullopt;
6509 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6510 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6511 if (auto Name = getModuleFlagNameSafely(*Flag))
6512 return *Name ==
"Debug Info Version";
6515 if (OpIt != ModFlags->op_end()) {
6516 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6523 bool BrokenDebugInfo =
false;
6526 if (!BrokenDebugInfo)
6532 M.getContext().diagnose(Diag);
6539 M.getContext().diagnose(DiagVersion);
6549 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6552 if (
F->hasFnAttribute(Attr)) {
6555 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6557 auto [Part, Rest] = S.
split(
',');
6563 const unsigned Dim = DimC -
'x';
6564 assert(Dim < 3 &&
"Unexpected dim char");
6574 F->addFnAttr(Attr, NewAttr);
6578 return S ==
"x" || S ==
"y" || S ==
"z";
6583 if (K ==
"kernel") {
6595 const unsigned Idx = (AlignIdxValuePair >> 16);
6596 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6601 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6606 if (K ==
"minctasm") {
6611 if (K ==
"maxnreg") {
6616 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6620 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6624 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6628 if (K ==
"grid_constant") {
6643 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6650 if (!SeenNodes.
insert(MD).second)
6657 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6664 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6666 const MDOperand &V = MD->getOperand(j + 1);
6669 NewOperands.
append({K, V});
6672 if (NewOperands.
size() > 1)
6685 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6686 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6687 if (ModRetainReleaseMarker) {
6693 ID->getString().split(ValueComp,
"#");
6694 if (ValueComp.
size() == 2) {
6695 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6699 M.eraseNamedMetadata(ModRetainReleaseMarker);
6710 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6736 bool InvalidCast =
false;
6738 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6751 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6753 Args.push_back(Arg);
6760 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6765 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6778 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6786 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6787 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6788 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6789 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6790 {
"objc_autoreleaseReturnValue",
6791 llvm::Intrinsic::objc_autoreleaseReturnValue},
6792 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6793 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6794 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6795 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6796 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6797 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6798 {
"objc_release", llvm::Intrinsic::objc_release},
6799 {
"objc_retain", llvm::Intrinsic::objc_retain},
6800 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6801 {
"objc_retainAutoreleaseReturnValue",
6802 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6803 {
"objc_retainAutoreleasedReturnValue",
6804 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6805 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6806 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6807 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6808 {
"objc_unsafeClaimAutoreleasedReturnValue",
6809 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6810 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6811 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6812 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6813 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6814 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6815 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6816 {
"objc_arc_annotation_topdown_bbstart",
6817 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6818 {
"objc_arc_annotation_topdown_bbend",
6819 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6820 {
"objc_arc_annotation_bottomup_bbstart",
6821 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6822 {
"objc_arc_annotation_bottomup_bbend",
6823 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6825 for (
auto &
I : RuntimeFuncs)
6826 UpgradeToIntrinsic(
I.first,
I.second);
6850 std::optional<bool> UseAddressDisc;
6853 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6854 for (
const MDNode *Flag : ModFlags->operands()) {
6856 if (Name && (*Name ==
"ptrauth-init-fini" ||
6857 *Name ==
"ptrauth-init-fini-address-discrimination"))
6862 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6863 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6864 constexpr unsigned ExpectedAddressMarker = 1;
6867 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6870 bool HasAddressDisc;
6871 if (!CPA->hasAddressDiscriminator())
6872 HasAddressDisc =
false;
6873 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6874 HasAddressDisc =
true;
6878 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6881 UseAddressDisc = HasAddressDisc;
6882 return CPA->getPointer();
6886 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6889 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6891 if (!GV || !GV->hasInitializer())
6895 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6898 std::vector<Constant *> NewStructors;
6899 NewStructors.reserve(OldStructorsArray->getNumOperands());
6901 for (
Use &U : OldStructorsArray->operands()) {
6910 Func = UpgradeSinglePointer(Func);
6914 NewStructors.push_back(
6923 if (GlobalArraysToUpgrade.
empty())
6925 assert(UseAddressDisc.has_value());
6927 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6928 GV->setInitializer(NewInit);
6931 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6941 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6945 bool HasObjCFlag =
false, HasClassProperties =
false;
6946 bool HasSwiftVersionFlag =
false;
6947 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6954 if (
Op->getNumOperands() != 3)
6968 if (ID->getString() ==
"Objective-C Image Info Version")
6970 if (ID->getString() ==
"Objective-C Class Properties")
6971 HasClassProperties =
true;
6973 if (ID->getString() ==
"PIC Level") {
6974 if (
auto *Behavior =
6976 uint64_t V = Behavior->getLimitedValue();
6982 if (ID->getString() ==
"PIE Level")
6983 if (
auto *Behavior =
6990 if (ID->getString() ==
"branch-target-enforcement" ||
6991 ID->getString().starts_with(
"sign-return-address")) {
6992 if (
auto *Behavior =
6998 Op->getOperand(1),
Op->getOperand(2)};
7008 if (ID->getString() ==
"Objective-C Image Info Section") {
7011 Value->getString().split(ValueComp,
" ");
7012 if (ValueComp.
size() != 1) {
7013 std::string NewValue;
7014 for (
auto &S : ValueComp)
7015 NewValue += S.str();
7026 if (ID->getString() ==
"Objective-C Garbage Collection") {
7029 assert(Md->getValue() &&
"Expected non-empty metadata");
7030 auto Type = Md->getValue()->getType();
7033 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7034 if ((Val & 0xff) != Val) {
7035 HasSwiftVersionFlag =
true;
7036 SwiftABIVersion = (Val & 0xff00) >> 8;
7037 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7038 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7049 if (ID->getString() ==
"amdgpu_code_object_version") {
7052 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7061 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7090 if (HasObjCFlag && !HasClassProperties) {
7096 if (HasSwiftVersionFlag) {
7100 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7102 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7110 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7114 auto MatchesVersion = [](
const MDNode *
Op) {
7115 return Op->getNumOperands() >= 3 &&
7129 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7130 assert(
Op->getNumOperands() >= 2 &&
7131 "Expected at least 2 operands - name and linkage type");
7143 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7154 auto TrimSpaces = [](
StringRef Section) -> std::string {
7156 Section.split(Components,
',');
7161 for (
auto Component : Components)
7162 OS <<
',' << Component.trim();
7167 for (
auto &GV : M.globals()) {
7168 if (!GV.hasSection())
7173 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7178 GV.setSection(TrimSpaces(Section));
7194struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7195 StrictFPUpgradeVisitor() =
default;
7198 if (!
Call.isStrictFP())
7204 Call.removeFnAttr(Attribute::StrictFP);
7205 Call.addFnAttr(Attribute::NoBuiltin);
7210struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7211 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7212 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7214 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7229 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7230 StrictFPUpgradeVisitor SFPV;
7235 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7236 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7237 for (
auto &Arg :
F.args())
7239 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7241 bool AddingAttrs =
false, RemovingAttrs =
false;
7242 AttrBuilder AttrsToAdd(
F.getContext());
7247 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7248 A.isValid() &&
A.isStringAttribute()) {
7249 F.setSection(
A.getValueAsString());
7251 RemovingAttrs =
true;
7255 A.isValid() &&
A.isStringAttribute()) {
7258 AddingAttrs = RemovingAttrs =
true;
7261 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7262 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7264 RemovingAttrs =
true;
7265 if (
A.getValueAsString() ==
"true") {
7266 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7275 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7278 if (
A.getValueAsBool()) {
7279 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7285 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7286 RemovingAttrs =
true;
7293 bool HandleDenormalMode =
false;
7295 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7298 DenormalFPMath = ParsedMode;
7300 AddingAttrs = RemovingAttrs =
true;
7301 HandleDenormalMode =
true;
7305 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7309 DenormalFPMathF32 = ParsedMode;
7311 AddingAttrs = RemovingAttrs =
true;
7312 HandleDenormalMode =
true;
7316 if (HandleDenormalMode)
7317 AttrsToAdd.addDenormalFPEnvAttr(
7321 F.removeFnAttrs(AttrsToRemove);
7324 F.addFnAttrs(AttrsToAdd);
7330 if (!
F.hasFnAttribute(FnAttrName))
7331 F.addFnAttr(FnAttrName,
Value);
7338 if (!
F.hasFnAttribute(FnAttrName)) {
7340 F.addFnAttr(FnAttrName);
7342 auto A =
F.getFnAttribute(FnAttrName);
7343 if (
"false" ==
A.getValueAsString())
7344 F.removeFnAttr(FnAttrName);
7345 else if (
"true" ==
A.getValueAsString()) {
7346 F.removeFnAttr(FnAttrName);
7347 F.addFnAttr(FnAttrName);
7353 Triple T(M.getTargetTriple());
7354 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7357 uint64_t BTEValue = 0;
7358 uint64_t BPPLRValue = 0;
7359 uint64_t GCSValue = 0;
7360 uint64_t SRAValue = 0;
7361 uint64_t SRAALLValue = 0;
7362 uint64_t SRABKeyValue = 0;
7364 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7368 if (
Op->getNumOperands() != 3)
7377 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7378 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7379 : IDStr ==
"guarded-control-stack" ? &GCSValue
7380 : IDStr ==
"sign-return-address" ? &SRAValue
7381 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7382 : IDStr ==
"sign-return-address-with-bkey"
7388 *ValPtr = CI->getZExtValue();
7394 bool BTE = BTEValue == 1;
7395 bool BPPLR = BPPLRValue == 1;
7396 bool GCS = GCSValue == 1;
7397 bool SRA = SRAValue == 1;
7400 if (SRA && SRAALLValue == 1)
7401 SignTypeValue =
"all";
7404 if (SRA && SRABKeyValue == 1)
7405 SignKeyValue =
"b_key";
7407 for (
Function &
F : M.getFunctionList()) {
7408 if (
F.isDeclaration())
7415 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7416 A.isValid() &&
"none" ==
A.getValueAsString()) {
7417 F.removeFnAttr(
"sign-return-address");
7418 F.removeFnAttr(
"sign-return-address-key");
7434 if (SRAALLValue == 1)
7436 if (SRABKeyValue == 1)
7463 if (
T->getNumOperands() < 1)
7468 if (S->getString().starts_with(
"llvm.vectorizer."))
7474 StringRef OldPrefix =
"llvm.vectorizer.";
7477 if (OldTag ==
"llvm.vectorizer.unroll")
7489 if (
T->getNumOperands() < 1)
7501 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7514 Ops.reserve(
T->getNumOperands());
7515 Ops.push_back(NewTag);
7516 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7517 Ops.push_back(
T->getOperand(
I));
7534 if (
T->isDistinct()) {
7535 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7547 Ops.reserve(
T->getNumOperands());
7558 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7559 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7560 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7563 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7565 auto I =
DL.find(
"-n64-");
7567 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7572 std::string Res =
DL.str();
7575 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7576 Res.append(Res.empty() ?
"G1" :
"-G1");
7584 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7585 Res.append(
"-ni:7:8:9");
7587 if (
DL.ends_with(
"ni:7"))
7589 if (
DL.ends_with(
"ni:7:8"))
7594 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7595 Res.append(
"-p7:160:256:256:32");
7596 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7597 Res.append(
"-p8:128:128:128:48");
7598 constexpr StringRef OldP8(
"-p8:128:128-");
7599 if (
DL.contains(OldP8))
7600 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7601 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7602 Res.append(
"-p9:192:256:256:32");
7606 if (!
DL.contains(
"m:e"))
7607 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7612 if (
T.isSystemZ() && !
DL.empty()) {
7614 if (!
DL.contains(
"-S64"))
7615 return "E-S64" +
DL.drop_front(1).str();
7619 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7622 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7623 if (!
DL.contains(AddrSpaces)) {
7625 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7626 if (R.match(Res, &
Groups))
7632 if (
T.isAArch64()) {
7634 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7635 Res.append(
"-Fn32");
7636 AddPtr32Ptr64AddrSpaces();
7640 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7644 std::string I64 =
"-i64:64";
7645 std::string I128 =
"-i128:128";
7647 size_t Pos = Res.find(I64);
7648 if (Pos !=
size_t(-1))
7649 Res.insert(Pos + I64.size(), I128);
7653 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7654 size_t Pos = Res.find(
"-S128");
7657 Res.insert(Pos,
"-f64:32:64");
7663 AddPtr32Ptr64AddrSpaces();
7671 if (!
T.isOSIAMCU()) {
7672 std::string I128 =
"-i128:128";
7675 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7676 if (R.match(Res, &
Groups))
7684 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7686 auto I =
Ref.find(
"-f80:32-");
7688 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7696 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7699 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7700 B.removeAttribute(
"no-frame-pointer-elim");
7702 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7704 if (FramePointer !=
"all")
7705 FramePointer =
"non-leaf";
7706 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7708 if (!FramePointer.
empty())
7709 B.addAttribute(
"frame-pointer", FramePointer);
7711 A =
B.getAttribute(
"null-pointer-is-valid");
7714 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7715 B.removeAttribute(
"null-pointer-is-valid");
7716 if (NullPointerIsValid)
7717 B.addAttribute(Attribute::NullPointerIsValid);
7720 A =
B.getAttribute(
"uniform-work-group-size");
7724 bool IsTrue = Val ==
"true";
7725 B.removeAttribute(
"uniform-work-group-size");
7727 B.addAttribute(
"uniform-work-group-size");
7738 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.