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)
1394 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1399 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1402 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1405 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1406 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1414 if (!Name.consume_front(
"tcgen05.mma."))
1418 if (Name.starts_with(
"ws"))
1421 return F->getIntrinsicID();
1425 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1426 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1427 Name.consume_front(
"param");
1431 if (!Name.consume_front(
"vp."))
1460 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1461 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1468 if (!Name.consume_front(
"vp."))
1488 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1489 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1494 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1506 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1517 if (Name.starts_with(
"to.fp16")) {
1521 FuncTy->getReturnType());
1524 if (Name.starts_with(
"from.fp16")) {
1528 FuncTy->getReturnType());
1540 if (Defaults.empty())
1552 if (
F->arg_size() >= FullDecl->
arg_size())
1557 if (
F->arg_size() < FirstDefault)
1565 bool CanUpgradeDebugIntrinsicsToRecords) {
1566 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1571 if (!Name.consume_front(
"llvm.") || Name.empty())
1577 bool IsArm = Name.consume_front(
"arm.");
1578 if (IsArm || Name.consume_front(
"aarch64.")) {
1584 if (Name.consume_front(
"amdgcn.")) {
1585 if (Name ==
"alignbit") {
1588 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1592 if (Name.consume_front(
"atomic.")) {
1593 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1594 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1603 switch (
F->getIntrinsicID()) {
1607 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1608 if (
F->arg_size() == 7) {
1613 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1614 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1615 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1616 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1617 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1618 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1619 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1620 if (
F->arg_size() == 8) {
1627 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1628 Name.consume_front(
"flat.atomic.")) {
1629 if (Name.starts_with(
"fadd") ||
1631 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1632 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1640 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1645 if (Name.starts_with(
"ldexp.")) {
1648 F->getParent(), Intrinsic::ldexp,
1649 {F->getReturnType(), F->getArg(1)->getType()});
1658 if (
F->arg_size() == 1) {
1659 if (Name.consume_front(
"convert.")) {
1673 F->arg_begin()->getType());
1679 if (Name ==
"coro.end" &&
1680 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1681 CoroEndID = Intrinsic::coro_end;
1682 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1683 CoroEndID = Intrinsic::coro_end_async;
1694 if (Name.consume_front(
"dbg.")) {
1696 if (CanUpgradeDebugIntrinsicsToRecords) {
1697 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1698 Name ==
"declare" || Name ==
"label") {
1707 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1710 Intrinsic::dbg_value);
1717 if (Name.consume_front(
"experimental.vector.")) {
1723 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1724 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1725 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1726 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1727 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1729 Intrinsic::vector_partial_reduce_add)
1732 const auto *FT =
F->getFunctionType();
1734 if (ID == Intrinsic::vector_extract ||
1735 ID == Intrinsic::vector_interleave2)
1738 if (ID != Intrinsic::vector_interleave2)
1740 if (ID == Intrinsic::vector_insert ||
1741 ID == Intrinsic::vector_partial_reduce_add)
1749 if (Name.consume_front(
"reduce.")) {
1751 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1752 if (R.match(Name, &
Groups))
1754 .
Case(
"add", Intrinsic::vector_reduce_add)
1755 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1756 .
Case(
"and", Intrinsic::vector_reduce_and)
1757 .
Case(
"or", Intrinsic::vector_reduce_or)
1758 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1759 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1760 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1761 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1762 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1763 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1764 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1769 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1774 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1775 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1780 auto Args =
F->getFunctionType()->params();
1782 {Args[V2 ? 1 : 0]});
1788 if (Name.consume_front(
"splice"))
1792 if (Name.consume_front(
"experimental.stepvector.")) {
1796 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1801 if (Name.starts_with(
"flt.rounds")) {
1804 Intrinsic::get_rounding);
1809 if (Name.starts_with(
"invariant.group.barrier")) {
1811 auto Args =
F->getFunctionType()->params();
1812 Type* ObjectPtr[1] = {Args[0]};
1815 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1820 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1821 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1822 if (IsLifetimeStart || IsLifetimeEnd) {
1823 if (
F->arg_size() == 2) {
1824 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1825 : Intrinsic::lifetime_end;
1830 F->getArg(1)->getType());
1832 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1852 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1853 .StartsWith(
"memmove.", Intrinsic::memmove)
1855 if (
F->arg_size() == 5) {
1859 F->getFunctionType()->params().slice(0, 3);
1865 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1868 const auto *FT =
F->getFunctionType();
1869 Type *ParamTypes[2] = {
1870 FT->getParamType(0),
1874 Intrinsic::memset, ParamTypes);
1880 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1881 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1882 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1883 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1885 if (MaskedID &&
F->arg_size() == 4) {
1887 if (MaskedID == Intrinsic::masked_load ||
1888 MaskedID == Intrinsic::masked_gather) {
1890 F->getParent(), MaskedID,
1891 {F->getReturnType(), F->getArg(0)->getType()});
1895 F->getParent(), MaskedID,
1896 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1902 if (Name.consume_front(
"nvvm.")) {
1904 if (
F->arg_size() == 1) {
1907 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1908 .Case(
"clz.i", Intrinsic::ctlz)
1909 .
Case(
"popc.i", Intrinsic::ctpop)
1913 {F->getReturnType()});
1916 }
else if (
F->arg_size() == 2) {
1919 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1920 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1921 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1922 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1926 {F->getReturnType()});
1963 F->getParent(), IID,
F->getReturnType(),
1964 F->getFunctionType()->params());
1975 {F->getArg(0)->getType()});
2000 bool Expand =
false;
2001 if (Name.consume_front(
"abs."))
2004 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2005 else if (Name.consume_front(
"fabs."))
2007 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2008 else if (Name.consume_front(
"ex2.approx."))
2011 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2012 else if (Name.consume_front(
"atomic.load."))
2021 else if (Name.consume_front(
"atomic."))
2036 else if (Name.consume_front(
"bitcast."))
2039 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2040 else if (Name.consume_front(
"rotate."))
2042 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2043 else if (Name.consume_front(
"ptr.gen.to."))
2046 else if (Name.consume_front(
"ptr."))
2049 else if (Name.consume_front(
"ldg.global."))
2051 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2052 Name.starts_with(
"p."));
2055 .
Case(
"barrier0",
true)
2056 .
Case(
"barrier.n",
true)
2057 .
Case(
"barrier.sync.cnt",
true)
2058 .
Case(
"barrier.sync",
true)
2059 .
Case(
"barrier",
true)
2060 .
Case(
"bar.sync",
true)
2061 .
Case(
"barrier0.popc",
true)
2062 .
Case(
"barrier0.and",
true)
2063 .
Case(
"barrier0.or",
true)
2064 .
Case(
"clz.ll",
true)
2065 .
Case(
"popc.ll",
true)
2067 .
Case(
"swap.lo.hi.b64",
true)
2068 .
Case(
"tanh.approx.f32",
true)
2080 if (Name.starts_with(
"objectsize.")) {
2081 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2082 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2085 Intrinsic::objectsize, Tys);
2092 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2095 F->getParent(), Intrinsic::ptr_annotation,
2096 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2102 if (Name.consume_front(
"riscv.")) {
2105 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2106 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2107 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2108 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2111 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2124 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2125 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2134 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2135 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2136 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2137 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2142 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2151 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2153 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2162 if (Name ==
"stackprotectorcheck") {
2169 if (Name ==
"thread.pointer") {
2171 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2177 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2180 F->getParent(), Intrinsic::var_annotation,
2181 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2184 if (Name.consume_front(
"vector.splice")) {
2185 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2195 if (Name.consume_front(
"wasm.")) {
2198 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2199 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2200 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2205 F->getReturnType());
2209 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2211 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2213 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2232 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2242 std::string
Name =
F->getName().str();
2245 Name,
F->getParent());
2256 if (Result != std::nullopt) {
2272 bool CanUpgradeDebugIntrinsicsToRecords) {
2292 GV->
getName() ==
"llvm.global_dtors")) ||
2307 unsigned N =
Init->getNumOperands();
2308 std::vector<Constant *> NewCtors(
N);
2309 for (
unsigned i = 0; i !=
N; ++i) {
2312 Ctor->getAggregateElement(1),
2326 unsigned NumElts = ResultTy->getNumElements() * 8;
2330 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2340 for (
unsigned l = 0; l != NumElts; l += 16)
2341 for (
unsigned i = 0; i != 16; ++i) {
2342 unsigned Idx = NumElts + i - Shift;
2344 Idx -= NumElts - 16;
2345 Idxs[l + i] = Idx + l;
2348 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2352 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2360 unsigned NumElts = ResultTy->getNumElements() * 8;
2364 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2374 for (
unsigned l = 0; l != NumElts; l += 16)
2375 for (
unsigned i = 0; i != 16; ++i) {
2376 unsigned Idx = i + Shift;
2378 Idx += NumElts - 16;
2379 Idxs[l + i] = Idx + l;
2382 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2386 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2394 Mask = Builder.CreateBitCast(Mask, MaskTy);
2400 for (
unsigned i = 0; i != NumElts; ++i)
2402 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2413 if (
C->isAllOnesValue())
2418 return Builder.CreateSelect(Mask, Op0, Op1);
2425 if (
C->isAllOnesValue())
2429 Mask->getType()->getIntegerBitWidth());
2430 Mask = Builder.CreateBitCast(Mask, MaskTy);
2431 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2432 return Builder.CreateSelect(Mask, Op0, Op1);
2445 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2446 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2451 ShiftVal &= (NumElts - 1);
2460 if (ShiftVal > 16) {
2468 for (
unsigned l = 0; l < NumElts; l += 16) {
2469 for (
unsigned i = 0; i != 16; ++i) {
2470 unsigned Idx = ShiftVal + i;
2471 if (!IsVALIGN && Idx >= 16)
2472 Idx += NumElts - 16;
2473 Indices[l + i] = Idx + l;
2478 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2484 bool ZeroMask,
bool IndexForm) {
2487 unsigned EltWidth = Ty->getScalarSizeInBits();
2488 bool IsFloat = Ty->isFPOrFPVectorTy();
2490 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2491 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2492 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2493 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2494 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2495 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2496 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2497 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2498 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2499 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2500 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2501 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2502 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2503 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2504 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2505 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2506 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2507 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2508 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2509 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2510 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2511 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2512 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2513 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2514 else if (VecWidth == 128 && EltWidth == 16)
2515 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2516 else if (VecWidth == 256 && EltWidth == 16)
2517 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2518 else if (VecWidth == 512 && EltWidth == 16)
2519 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2520 else if (VecWidth == 128 && EltWidth == 8)
2521 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2522 else if (VecWidth == 256 && EltWidth == 8)
2523 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2524 else if (VecWidth == 512 && EltWidth == 8)
2525 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2536 Value *V = Builder.CreateIntrinsic(IID, Args);
2548 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2559 bool IsRotateRight) {
2569 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2570 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2573 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2574 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2619 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2624 bool IsShiftRight,
bool ZeroMask) {
2638 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2639 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2642 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2643 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2658 const Align Alignment =
2660 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2665 if (
C->isAllOnesValue())
2666 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2671 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2677 const Align Alignment =
2686 if (
C->isAllOnesValue())
2687 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2692 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2698 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2699 {Op0, Builder.getInt1(
false)});
2714 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2715 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2716 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2717 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2718 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2721 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2722 LHS = Builder.CreateAnd(
LHS, Mask);
2723 RHS = Builder.CreateAnd(
RHS, Mask);
2740 if (!
C || !
C->isAllOnesValue())
2741 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2746 for (
unsigned i = 0; i != NumElts; ++i)
2748 for (
unsigned i = NumElts; i != 8; ++i)
2749 Indices[i] = NumElts + i % NumElts;
2750 Vec = Builder.CreateShuffleVector(Vec,
2754 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2758 unsigned CC,
bool Signed) {
2766 }
else if (CC == 7) {
2802 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2803 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2805 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2806 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2815 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2821 Name = Name.substr(12);
2826 if (Name.starts_with(
"max.p")) {
2827 if (VecWidth == 128 && EltWidth == 32)
2828 IID = Intrinsic::x86_sse_max_ps;
2829 else if (VecWidth == 128 && EltWidth == 64)
2830 IID = Intrinsic::x86_sse2_max_pd;
2831 else if (VecWidth == 256 && EltWidth == 32)
2832 IID = Intrinsic::x86_avx_max_ps_256;
2833 else if (VecWidth == 256 && EltWidth == 64)
2834 IID = Intrinsic::x86_avx_max_pd_256;
2837 }
else if (Name.starts_with(
"min.p")) {
2838 if (VecWidth == 128 && EltWidth == 32)
2839 IID = Intrinsic::x86_sse_min_ps;
2840 else if (VecWidth == 128 && EltWidth == 64)
2841 IID = Intrinsic::x86_sse2_min_pd;
2842 else if (VecWidth == 256 && EltWidth == 32)
2843 IID = Intrinsic::x86_avx_min_ps_256;
2844 else if (VecWidth == 256 && EltWidth == 64)
2845 IID = Intrinsic::x86_avx_min_pd_256;
2848 }
else if (Name.starts_with(
"pshuf.b.")) {
2849 if (VecWidth == 128)
2850 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2851 else if (VecWidth == 256)
2852 IID = Intrinsic::x86_avx2_pshuf_b;
2853 else if (VecWidth == 512)
2854 IID = Intrinsic::x86_avx512_pshuf_b_512;
2857 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2858 if (VecWidth == 128)
2859 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2860 else if (VecWidth == 256)
2861 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2862 else if (VecWidth == 512)
2863 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2866 }
else if (Name.starts_with(
"pmulh.w.")) {
2867 if (VecWidth == 128)
2868 IID = Intrinsic::x86_sse2_pmulh_w;
2869 else if (VecWidth == 256)
2870 IID = Intrinsic::x86_avx2_pmulh_w;
2871 else if (VecWidth == 512)
2872 IID = Intrinsic::x86_avx512_pmulh_w_512;
2875 }
else if (Name.starts_with(
"pmulhu.w.")) {
2876 if (VecWidth == 128)
2877 IID = Intrinsic::x86_sse2_pmulhu_w;
2878 else if (VecWidth == 256)
2879 IID = Intrinsic::x86_avx2_pmulhu_w;
2880 else if (VecWidth == 512)
2881 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2884 }
else if (Name.starts_with(
"pmaddw.d.")) {
2885 if (VecWidth == 128)
2886 IID = Intrinsic::x86_sse2_pmadd_wd;
2887 else if (VecWidth == 256)
2888 IID = Intrinsic::x86_avx2_pmadd_wd;
2889 else if (VecWidth == 512)
2890 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2893 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2894 if (VecWidth == 128)
2895 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2896 else if (VecWidth == 256)
2897 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2898 else if (VecWidth == 512)
2899 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2902 }
else if (Name.starts_with(
"packsswb.")) {
2903 if (VecWidth == 128)
2904 IID = Intrinsic::x86_sse2_packsswb_128;
2905 else if (VecWidth == 256)
2906 IID = Intrinsic::x86_avx2_packsswb;
2907 else if (VecWidth == 512)
2908 IID = Intrinsic::x86_avx512_packsswb_512;
2911 }
else if (Name.starts_with(
"packssdw.")) {
2912 if (VecWidth == 128)
2913 IID = Intrinsic::x86_sse2_packssdw_128;
2914 else if (VecWidth == 256)
2915 IID = Intrinsic::x86_avx2_packssdw;
2916 else if (VecWidth == 512)
2917 IID = Intrinsic::x86_avx512_packssdw_512;
2920 }
else if (Name.starts_with(
"packuswb.")) {
2921 if (VecWidth == 128)
2922 IID = Intrinsic::x86_sse2_packuswb_128;
2923 else if (VecWidth == 256)
2924 IID = Intrinsic::x86_avx2_packuswb;
2925 else if (VecWidth == 512)
2926 IID = Intrinsic::x86_avx512_packuswb_512;
2929 }
else if (Name.starts_with(
"packusdw.")) {
2930 if (VecWidth == 128)
2931 IID = Intrinsic::x86_sse41_packusdw;
2932 else if (VecWidth == 256)
2933 IID = Intrinsic::x86_avx2_packusdw;
2934 else if (VecWidth == 512)
2935 IID = Intrinsic::x86_avx512_packusdw_512;
2938 }
else if (Name.starts_with(
"vpermilvar.")) {
2939 if (VecWidth == 128 && EltWidth == 32)
2940 IID = Intrinsic::x86_avx_vpermilvar_ps;
2941 else if (VecWidth == 128 && EltWidth == 64)
2942 IID = Intrinsic::x86_avx_vpermilvar_pd;
2943 else if (VecWidth == 256 && EltWidth == 32)
2944 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2945 else if (VecWidth == 256 && EltWidth == 64)
2946 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2947 else if (VecWidth == 512 && EltWidth == 32)
2948 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2949 else if (VecWidth == 512 && EltWidth == 64)
2950 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2953 }
else if (Name ==
"cvtpd2dq.256") {
2954 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2955 }
else if (Name ==
"cvtpd2ps.256") {
2956 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2957 }
else if (Name ==
"cvttpd2dq.256") {
2958 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2959 }
else if (Name ==
"cvttps2dq.128") {
2960 IID = Intrinsic::x86_sse2_cvttps2dq;
2961 }
else if (Name ==
"cvttps2dq.256") {
2962 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2963 }
else if (Name.starts_with(
"permvar.")) {
2965 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2966 IID = Intrinsic::x86_avx2_permps;
2967 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2968 IID = Intrinsic::x86_avx2_permd;
2969 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2970 IID = Intrinsic::x86_avx512_permvar_df_256;
2971 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2972 IID = Intrinsic::x86_avx512_permvar_di_256;
2973 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2974 IID = Intrinsic::x86_avx512_permvar_sf_512;
2975 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2976 IID = Intrinsic::x86_avx512_permvar_si_512;
2977 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2978 IID = Intrinsic::x86_avx512_permvar_df_512;
2979 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2980 IID = Intrinsic::x86_avx512_permvar_di_512;
2981 else if (VecWidth == 128 && EltWidth == 16)
2982 IID = Intrinsic::x86_avx512_permvar_hi_128;
2983 else if (VecWidth == 256 && EltWidth == 16)
2984 IID = Intrinsic::x86_avx512_permvar_hi_256;
2985 else if (VecWidth == 512 && EltWidth == 16)
2986 IID = Intrinsic::x86_avx512_permvar_hi_512;
2987 else if (VecWidth == 128 && EltWidth == 8)
2988 IID = Intrinsic::x86_avx512_permvar_qi_128;
2989 else if (VecWidth == 256 && EltWidth == 8)
2990 IID = Intrinsic::x86_avx512_permvar_qi_256;
2991 else if (VecWidth == 512 && EltWidth == 8)
2992 IID = Intrinsic::x86_avx512_permvar_qi_512;
2995 }
else if (Name.starts_with(
"dbpsadbw.")) {
2996 if (VecWidth == 128)
2997 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2998 else if (VecWidth == 256)
2999 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3000 else if (VecWidth == 512)
3001 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3004 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3005 if (VecWidth == 128)
3006 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3007 else if (VecWidth == 256)
3008 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3009 else if (VecWidth == 512)
3010 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3013 }
else if (Name.starts_with(
"conflict.")) {
3014 if (Name[9] ==
'd' && VecWidth == 128)
3015 IID = Intrinsic::x86_avx512_conflict_d_128;
3016 else if (Name[9] ==
'd' && VecWidth == 256)
3017 IID = Intrinsic::x86_avx512_conflict_d_256;
3018 else if (Name[9] ==
'd' && VecWidth == 512)
3019 IID = Intrinsic::x86_avx512_conflict_d_512;
3020 else if (Name[9] ==
'q' && VecWidth == 128)
3021 IID = Intrinsic::x86_avx512_conflict_q_128;
3022 else if (Name[9] ==
'q' && VecWidth == 256)
3023 IID = Intrinsic::x86_avx512_conflict_q_256;
3024 else if (Name[9] ==
'q' && VecWidth == 512)
3025 IID = Intrinsic::x86_avx512_conflict_q_512;
3028 }
else if (Name.starts_with(
"pavg.")) {
3029 if (Name[5] ==
'b' && VecWidth == 128)
3030 IID = Intrinsic::x86_sse2_pavg_b;
3031 else if (Name[5] ==
'b' && VecWidth == 256)
3032 IID = Intrinsic::x86_avx2_pavg_b;
3033 else if (Name[5] ==
'b' && VecWidth == 512)
3034 IID = Intrinsic::x86_avx512_pavg_b_512;
3035 else if (Name[5] ==
'w' && VecWidth == 128)
3036 IID = Intrinsic::x86_sse2_pavg_w;
3037 else if (Name[5] ==
'w' && VecWidth == 256)
3038 IID = Intrinsic::x86_avx2_pavg_w;
3039 else if (Name[5] ==
'w' && VecWidth == 512)
3040 IID = Intrinsic::x86_avx512_pavg_w_512;
3049 Rep = Builder.CreateIntrinsic(IID, Args);
3060 if (AsmStr->find(
"mov\tfp") == 0 &&
3061 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3062 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3063 AsmStr->replace(Pos, 1,
";");
3069 Value *Rep =
nullptr;
3071 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3073 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3074 {Arg, Builder.getTrue()},
3076 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3077 Type *Ty = (Name ==
"abs.bf16")
3081 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3082 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3083 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3084 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3085 : Intrinsic::nvvm_fabs;
3086 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3087 }
else if (Name.consume_front(
"ex2.approx.")) {
3089 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3090 : Intrinsic::nvvm_ex2_approx;
3091 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3092 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3093 Name.starts_with(
"atomic.load.add.f64.p")) {
3096 Rep = Builder.CreateAtomicRMW(
3102 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3103 Name.starts_with(
"atomic.load.dec.32.p")) {
3108 Rep = Builder.CreateAtomicRMW(
3112 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3118 Op.contains(
".cta.") ?
"block" :
"");
3119 if (
Op.starts_with(
"cas.")) {
3121 Value *Pair = Builder.CreateAtomicCmpXchg(
3124 Rep = Builder.CreateExtractValue(Pair, 0);
3142 "unexpected nvvm scoped atomic intrinsic");
3143 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3146 }
else if (Name ==
"clz.ll") {
3149 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3150 {Arg, Builder.getFalse()},
3152 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3153 }
else if (Name ==
"popc.ll") {
3157 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3158 Arg,
nullptr,
"ctpop");
3159 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3160 }
else if (Name ==
"h2f") {
3162 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3163 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3164 }
else if (Name.consume_front(
"bitcast.") &&
3165 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3168 }
else if (Name ==
"rotate.b32") {
3171 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3172 {Arg, Arg, ShiftAmt});
3173 }
else if (Name ==
"rotate.b64") {
3177 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3178 {Arg, Arg, ZExtShiftAmt});
3179 }
else if (Name ==
"rotate.right.b64") {
3183 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3184 {Arg, Arg, ZExtShiftAmt});
3185 }
else if (Name ==
"swap.lo.hi.b64") {
3188 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3189 {Arg, Arg, Builder.getInt64(32)});
3190 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3193 Name.starts_with(
".to.gen"))) {
3195 }
else if (Name.consume_front(
"ldg.global")) {
3199 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3202 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3204 }
else if (Name ==
"tanh.approx.f32") {
3208 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3210 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3212 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3213 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3215 }
else if (Name ==
"barrier") {
3216 Rep = Builder.CreateIntrinsic(
3217 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3219 }
else if (Name ==
"barrier.sync") {
3220 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3222 }
else if (Name ==
"barrier.sync.cnt") {
3223 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3225 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3226 Name ==
"barrier0.or") {
3228 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3232 .
Case(
"barrier0.popc",
3233 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3234 .
Case(
"barrier0.and",
3235 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3236 .
Case(
"barrier0.or",
3237 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3238 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3239 Rep = Builder.CreateZExt(Bar, CI->
getType());
3253 ? Builder.CreateBitCast(Arg, NewType)
3256 Rep = Builder.CreateCall(NewFn, Args);
3257 if (
F->getReturnType()->isIntegerTy())
3258 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3268 Value *Rep =
nullptr;
3270 if (Name.starts_with(
"sse4a.movnt.")) {
3282 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3285 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3286 }
else if (Name.starts_with(
"avx.movnt.") ||
3287 Name.starts_with(
"avx512.storent.")) {
3299 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3300 }
else if (Name ==
"sse2.storel.dq") {
3305 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3306 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3307 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3308 }
else if (Name.starts_with(
"sse.storeu.") ||
3309 Name.starts_with(
"sse2.storeu.") ||
3310 Name.starts_with(
"avx.storeu.")) {
3313 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3314 }
else if (Name ==
"avx512.mask.store.ss") {
3318 }
else if (Name.starts_with(
"avx512.mask.store")) {
3320 bool Aligned = Name[17] !=
'u';
3323 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3326 bool CmpEq = Name[9] ==
'e';
3329 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3330 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3337 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3338 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3340 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3341 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3342 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3343 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3344 Name.starts_with(
"sse2.sqrt.p") ||
3345 Name.starts_with(
"sse.sqrt.p")) {
3346 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3347 {CI->getArgOperand(0)});
3348 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3352 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3353 : Intrinsic::x86_avx512_sqrt_pd_512;
3356 Rep = Builder.CreateIntrinsic(IID, Args);
3358 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3359 {CI->getArgOperand(0)});
3363 }
else if (Name.starts_with(
"avx512.ptestm") ||
3364 Name.starts_with(
"avx512.ptestnm")) {
3368 Rep = Builder.CreateAnd(Op0, Op1);
3374 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3376 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3379 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3382 }
else if (Name.starts_with(
"avx512.kunpck")) {
3387 for (
unsigned i = 0; i != NumElts; ++i)
3396 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3397 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3398 }
else if (Name ==
"avx512.kand.w") {
3401 Rep = Builder.CreateAnd(
LHS,
RHS);
3402 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3403 }
else if (Name ==
"avx512.kandn.w") {
3406 LHS = Builder.CreateNot(
LHS);
3407 Rep = Builder.CreateAnd(
LHS,
RHS);
3408 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3409 }
else if (Name ==
"avx512.kor.w") {
3412 Rep = Builder.CreateOr(
LHS,
RHS);
3413 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3414 }
else if (Name ==
"avx512.kxor.w") {
3417 Rep = Builder.CreateXor(
LHS,
RHS);
3418 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3419 }
else if (Name ==
"avx512.kxnor.w") {
3422 LHS = Builder.CreateNot(
LHS);
3423 Rep = Builder.CreateXor(
LHS,
RHS);
3424 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3425 }
else if (Name ==
"avx512.knot.w") {
3427 Rep = Builder.CreateNot(Rep);
3428 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3429 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3432 Rep = Builder.CreateOr(
LHS,
RHS);
3433 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3435 if (Name[14] ==
'c')
3439 Rep = Builder.CreateICmpEQ(Rep,
C);
3440 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3441 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3442 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3443 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3444 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3447 ConstantInt::get(I32Ty, 0));
3449 ConstantInt::get(I32Ty, 0));
3451 if (Name.contains(
".add."))
3452 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3453 else if (Name.contains(
".sub."))
3454 EltOp = Builder.CreateFSub(Elt0, Elt1);
3455 else if (Name.contains(
".mul."))
3456 EltOp = Builder.CreateFMul(Elt0, Elt1);
3458 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3459 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3460 ConstantInt::get(I32Ty, 0));
3461 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3463 bool CmpEq = Name[16] ==
'e';
3465 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3467 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3474 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3477 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3480 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3487 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3489 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3490 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3492 if (VecWidth == 128 && EltWidth == 32)
3493 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3494 else if (VecWidth == 256 && EltWidth == 32)
3495 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3496 else if (VecWidth == 512 && EltWidth == 32)
3497 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3498 else if (VecWidth == 128 && EltWidth == 64)
3499 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3500 else if (VecWidth == 256 && EltWidth == 64)
3501 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3502 else if (VecWidth == 512 && EltWidth == 64)
3503 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3510 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3513 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3514 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3516 if (VecWidth == 128 && EltWidth == 32)
3517 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3518 else if (VecWidth == 256 && EltWidth == 32)
3519 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3520 else if (VecWidth == 512 && EltWidth == 32)
3521 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3522 else if (VecWidth == 128 && EltWidth == 64)
3523 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3524 else if (VecWidth == 256 && EltWidth == 64)
3525 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3526 else if (VecWidth == 512 && EltWidth == 64)
3527 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3532 if (VecWidth == 512)
3534 Args.push_back(Mask);
3536 Rep = Builder.CreateIntrinsic(IID, Args);
3537 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3541 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3544 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3545 Name.starts_with(
"avx512.cvtw2mask.") ||
3546 Name.starts_with(
"avx512.cvtd2mask.") ||
3547 Name.starts_with(
"avx512.cvtq2mask.")) {
3552 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3553 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3554 Name.starts_with(
"avx512.mask.pabs")) {
3556 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3557 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3558 Name.starts_with(
"avx512.mask.pmaxs")) {
3560 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3561 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3562 Name.starts_with(
"avx512.mask.pmaxu")) {
3564 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3565 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3566 Name.starts_with(
"avx512.mask.pmins")) {
3568 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3569 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3570 Name.starts_with(
"avx512.mask.pminu")) {
3572 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3573 Name ==
"avx512.pmulu.dq.512" ||
3574 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3576 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3577 Name ==
"avx512.pmul.dq.512" ||
3578 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3580 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3581 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3586 }
else if (Name ==
"avx512.cvtusi2sd") {
3591 }
else if (Name ==
"sse2.cvtss2sd") {
3593 Rep = Builder.CreateFPExt(
3596 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3597 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3598 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3599 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3600 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3601 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3602 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3603 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3604 Name ==
"avx512.mask.cvtqq2ps.256" ||
3605 Name ==
"avx512.mask.cvtqq2ps.512" ||
3606 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3607 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3608 Name ==
"avx.cvt.ps2.pd.256" ||
3609 Name ==
"avx512.mask.cvtps2pd.128" ||
3610 Name ==
"avx512.mask.cvtps2pd.256") {
3615 unsigned NumDstElts = DstTy->getNumElements();
3616 if (NumDstElts < SrcTy->getNumElements()) {
3617 assert(NumDstElts == 2 &&
"Unexpected vector size");
3618 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3621 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3622 bool IsUnsigned = Name.contains(
"cvtu");
3624 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3628 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3629 : Intrinsic::x86_avx512_sitofp_round;
3630 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3633 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3634 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3640 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3641 Name.starts_with(
"vcvtph2ps.")) {
3645 unsigned NumDstElts = DstTy->getNumElements();
3646 if (NumDstElts != SrcTy->getNumElements()) {
3647 assert(NumDstElts == 4 &&
"Unexpected vector size");
3648 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3650 Rep = Builder.CreateBitCast(
3652 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3656 }
else if (Name.starts_with(
"avx512.mask.load")) {
3658 bool Aligned = Name[16] !=
'u';
3661 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3665 ResultTy->getNumElements());
3666 Rep = Builder.CreateIntrinsic(
3667 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3669 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3675 Rep = Builder.CreateIntrinsic(
3676 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3678 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3679 Name.starts_with(
"avx512.mask.expand.")) {
3683 ResultTy->getNumElements());
3685 bool IsCompress = Name[12] ==
'c';
3686 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3687 : Intrinsic::x86_avx512_mask_expand;
3688 Rep = Builder.CreateIntrinsic(
3690 }
else if (Name.starts_with(
"xop.vpcom")) {
3692 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3693 Name.ends_with(
"uq"))
3695 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3696 Name.ends_with(
"d") || Name.ends_with(
"q"))
3705 Name = Name.substr(9);
3706 if (Name.starts_with(
"lt"))
3708 else if (Name.starts_with(
"le"))
3710 else if (Name.starts_with(
"gt"))
3712 else if (Name.starts_with(
"ge"))
3714 else if (Name.starts_with(
"eq"))
3716 else if (Name.starts_with(
"ne"))
3718 else if (Name.starts_with(
"false"))
3720 else if (Name.starts_with(
"true"))
3727 }
else if (Name.starts_with(
"xop.vpcmov")) {
3729 Value *NotSel = Builder.CreateNot(Sel);
3732 Rep = Builder.CreateOr(Sel0, Sel1);
3733 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3734 Name.starts_with(
"avx512.mask.prol")) {
3736 }
else if (Name.starts_with(
"avx512.pror") ||
3737 Name.starts_with(
"avx512.mask.pror")) {
3739 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3740 Name.starts_with(
"avx512.mask.vpshld") ||
3741 Name.starts_with(
"avx512.maskz.vpshld")) {
3742 bool ZeroMask = Name[11] ==
'z';
3744 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3745 Name.starts_with(
"avx512.mask.vpshrd") ||
3746 Name.starts_with(
"avx512.maskz.vpshrd")) {
3747 bool ZeroMask = Name[11] ==
'z';
3749 }
else if (Name ==
"sse42.crc32.64.8") {
3752 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3754 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3755 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3756 Name.starts_with(
"avx512.vbroadcast.s")) {
3759 Type *EltTy = VecTy->getElementType();
3760 unsigned EltNum = VecTy->getNumElements();
3764 for (
unsigned I = 0;
I < EltNum; ++
I)
3765 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3766 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3767 Name.starts_with(
"sse41.pmovzx") ||
3768 Name.starts_with(
"avx2.pmovsx") ||
3769 Name.starts_with(
"avx2.pmovzx") ||
3770 Name.starts_with(
"avx512.mask.pmovsx") ||
3771 Name.starts_with(
"avx512.mask.pmovzx")) {
3773 unsigned NumDstElts = DstTy->getNumElements();
3777 for (
unsigned i = 0; i != NumDstElts; ++i)
3782 bool DoSext = Name.contains(
"pmovsx");
3784 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3789 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3790 Name ==
"avx512.mask.pmov.qd.512" ||
3791 Name ==
"avx512.mask.pmov.wb.256" ||
3792 Name ==
"avx512.mask.pmov.wb.512") {
3797 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3798 Name ==
"avx2.vbroadcasti128") {
3804 if (NumSrcElts == 2)
3807 Rep = Builder.CreateShuffleVector(
Load,
3809 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3810 Name.starts_with(
"avx512.mask.shuf.f")) {
3815 unsigned ControlBitsMask = NumLanes - 1;
3816 unsigned NumControlBits = NumLanes / 2;
3819 for (
unsigned l = 0; l != NumLanes; ++l) {
3820 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3822 if (l >= NumLanes / 2)
3823 LaneMask += NumLanes;
3824 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3825 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3831 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3832 Name.starts_with(
"avx512.mask.broadcasti")) {
3835 unsigned NumDstElts =
3839 for (
unsigned i = 0; i != NumDstElts; ++i)
3840 ShuffleMask[i] = i % NumSrcElts;
3846 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3847 Name.starts_with(
"avx2.vbroadcast") ||
3848 Name.starts_with(
"avx512.pbroadcast") ||
3849 Name.starts_with(
"avx512.mask.broadcast.s")) {
3856 Rep = Builder.CreateShuffleVector(
Op, M);
3861 }
else if (Name.starts_with(
"sse2.padds.") ||
3862 Name.starts_with(
"avx2.padds.") ||
3863 Name.starts_with(
"avx512.padds.") ||
3864 Name.starts_with(
"avx512.mask.padds.")) {
3866 }
else if (Name.starts_with(
"sse2.psubs.") ||
3867 Name.starts_with(
"avx2.psubs.") ||
3868 Name.starts_with(
"avx512.psubs.") ||
3869 Name.starts_with(
"avx512.mask.psubs.")) {
3871 }
else if (Name.starts_with(
"sse2.paddus.") ||
3872 Name.starts_with(
"avx2.paddus.") ||
3873 Name.starts_with(
"avx512.mask.paddus.")) {
3875 }
else if (Name.starts_with(
"sse2.psubus.") ||
3876 Name.starts_with(
"avx2.psubus.") ||
3877 Name.starts_with(
"avx512.mask.psubus.")) {
3879 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3884 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3888 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3893 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3898 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3899 Name ==
"avx512.psll.dq.512") {
3903 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3904 Name ==
"avx512.psrl.dq.512") {
3908 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3909 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3910 Name.starts_with(
"avx2.pblendd.")) {
3915 unsigned NumElts = VecTy->getNumElements();
3918 for (
unsigned i = 0; i != NumElts; ++i)
3919 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
3921 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3922 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3923 Name ==
"avx2.vinserti128" ||
3924 Name.starts_with(
"avx512.mask.insert")) {
3928 unsigned DstNumElts =
3930 unsigned SrcNumElts =
3932 unsigned Scale = DstNumElts / SrcNumElts;
3939 for (
unsigned i = 0; i != SrcNumElts; ++i)
3941 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3942 Idxs[i] = SrcNumElts;
3943 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3957 for (
unsigned i = 0; i != DstNumElts; ++i)
3960 for (
unsigned i = 0; i != SrcNumElts; ++i)
3961 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
3962 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3968 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3969 Name ==
"avx2.vextracti128" ||
3970 Name.starts_with(
"avx512.mask.vextract")) {
3973 unsigned DstNumElts =
3975 unsigned SrcNumElts =
3977 unsigned Scale = SrcNumElts / DstNumElts;
3984 for (
unsigned i = 0; i != DstNumElts; ++i) {
3985 Idxs[i] = i + (
Imm * DstNumElts);
3987 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3993 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3994 Name.starts_with(
"avx512.mask.perm.di.")) {
3998 unsigned NumElts = VecTy->getNumElements();
4001 for (
unsigned i = 0; i != NumElts; ++i)
4002 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4004 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4009 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4021 unsigned HalfSize = NumElts / 2;
4033 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4034 for (
unsigned i = 0; i < HalfSize; ++i)
4035 ShuffleMask[i] = StartIndex + i;
4038 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4039 for (
unsigned i = 0; i < HalfSize; ++i)
4040 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4042 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4044 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4045 Name.starts_with(
"avx512.mask.vpermil.p") ||
4046 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4050 unsigned NumElts = VecTy->getNumElements();
4052 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4053 unsigned IdxMask = ((1 << IdxSize) - 1);
4059 for (
unsigned i = 0; i != NumElts; ++i)
4060 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4062 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4067 }
else if (Name ==
"sse2.pshufl.w" ||
4068 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4073 if (Name ==
"sse2.pshufl.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] = ((
Imm >> (2 * i)) & 0x3) + l;
4080 for (
unsigned i = 4; i != 8; ++i)
4081 Idxs[i + l] = i + l;
4084 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4089 }
else if (Name ==
"sse2.pshufh.w" ||
4090 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4095 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4099 for (
unsigned l = 0; l != NumElts; l += 8) {
4100 for (
unsigned i = 0; i != 4; ++i)
4101 Idxs[i + l] = i + l;
4102 for (
unsigned i = 0; i != 4; ++i)
4103 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4106 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4111 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4118 unsigned HalfLaneElts = NumLaneElts / 2;
4121 for (
unsigned i = 0; i != NumElts; ++i) {
4123 Idxs[i] = i - (i % NumLaneElts);
4125 if ((i % NumLaneElts) >= HalfLaneElts)
4129 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4132 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4136 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4137 Name.starts_with(
"avx512.mask.movshdup") ||
4138 Name.starts_with(
"avx512.mask.movsldup")) {
4144 if (Name.starts_with(
"avx512.mask.movshdup."))
4148 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4149 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4150 Idxs[i + l + 0] = i + l +
Offset;
4151 Idxs[i + l + 1] = i + l +
Offset;
4154 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4158 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4159 Name.starts_with(
"avx512.mask.unpckl.")) {
4166 for (
int l = 0; l != NumElts; l += NumLaneElts)
4167 for (
int i = 0; i != NumLaneElts; ++i)
4168 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4170 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4174 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4175 Name.starts_with(
"avx512.mask.unpckh.")) {
4182 for (
int l = 0; l != NumElts; l += NumLaneElts)
4183 for (
int i = 0; i != NumLaneElts; ++i)
4184 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4186 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4190 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4191 Name.starts_with(
"avx512.mask.pand.")) {
4194 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4196 Rep = Builder.CreateBitCast(Rep, FTy);
4199 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4200 Name.starts_with(
"avx512.mask.pandn.")) {
4203 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4204 Rep = Builder.CreateAnd(Rep,
4206 Rep = Builder.CreateBitCast(Rep, FTy);
4209 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4210 Name.starts_with(
"avx512.mask.por.")) {
4213 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4215 Rep = Builder.CreateBitCast(Rep, FTy);
4218 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4219 Name.starts_with(
"avx512.mask.pxor.")) {
4222 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4224 Rep = Builder.CreateBitCast(Rep, FTy);
4227 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4231 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4235 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4239 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4240 if (Name.ends_with(
".512")) {
4242 if (Name[17] ==
's')
4243 IID = Intrinsic::x86_avx512_add_ps_512;
4245 IID = Intrinsic::x86_avx512_add_pd_512;
4247 Rep = Builder.CreateIntrinsic(
4255 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4256 if (Name.ends_with(
".512")) {
4258 if (Name[17] ==
's')
4259 IID = Intrinsic::x86_avx512_div_ps_512;
4261 IID = Intrinsic::x86_avx512_div_pd_512;
4263 Rep = Builder.CreateIntrinsic(
4271 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4272 if (Name.ends_with(
".512")) {
4274 if (Name[17] ==
's')
4275 IID = Intrinsic::x86_avx512_mul_ps_512;
4277 IID = Intrinsic::x86_avx512_mul_pd_512;
4279 Rep = Builder.CreateIntrinsic(
4287 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4288 if (Name.ends_with(
".512")) {
4290 if (Name[17] ==
's')
4291 IID = Intrinsic::x86_avx512_sub_ps_512;
4293 IID = Intrinsic::x86_avx512_sub_pd_512;
4295 Rep = Builder.CreateIntrinsic(
4303 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4304 Name.starts_with(
"avx512.mask.min.p")) &&
4305 Name.drop_front(18) ==
".512") {
4306 bool IsDouble = Name[17] ==
'd';
4307 bool IsMin = Name[13] ==
'i';
4309 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4310 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4313 Rep = Builder.CreateIntrinsic(
4318 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4320 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4321 {CI->getArgOperand(0), Builder.getInt1(false)});
4324 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4325 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4326 bool IsVariable = Name[16] ==
'v';
4327 char Size = Name[16] ==
'.' ? Name[17]
4328 : Name[17] ==
'.' ? Name[18]
4329 : Name[18] ==
'.' ? Name[19]
4333 if (IsVariable && Name[17] !=
'.') {
4334 if (
Size ==
'd' && Name[17] ==
'2')
4335 IID = Intrinsic::x86_avx2_psllv_q;
4336 else if (
Size ==
'd' && Name[17] ==
'4')
4337 IID = Intrinsic::x86_avx2_psllv_q_256;
4338 else if (
Size ==
's' && Name[17] ==
'4')
4339 IID = Intrinsic::x86_avx2_psllv_d;
4340 else if (
Size ==
's' && Name[17] ==
'8')
4341 IID = Intrinsic::x86_avx2_psllv_d_256;
4342 else if (
Size ==
'h' && Name[17] ==
'8')
4343 IID = Intrinsic::x86_avx512_psllv_w_128;
4344 else if (
Size ==
'h' && Name[17] ==
'1')
4345 IID = Intrinsic::x86_avx512_psllv_w_256;
4346 else if (Name[17] ==
'3' && Name[18] ==
'2')
4347 IID = Intrinsic::x86_avx512_psllv_w_512;
4350 }
else if (Name.ends_with(
".128")) {
4352 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4353 : Intrinsic::x86_sse2_psll_d;
4354 else if (
Size ==
'q')
4355 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4356 : Intrinsic::x86_sse2_psll_q;
4357 else if (
Size ==
'w')
4358 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4359 : Intrinsic::x86_sse2_psll_w;
4362 }
else if (Name.ends_with(
".256")) {
4364 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4365 : Intrinsic::x86_avx2_psll_d;
4366 else if (
Size ==
'q')
4367 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4368 : Intrinsic::x86_avx2_psll_q;
4369 else if (
Size ==
'w')
4370 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4371 : Intrinsic::x86_avx2_psll_w;
4376 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4377 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4378 : Intrinsic::x86_avx512_psll_d_512;
4379 else if (
Size ==
'q')
4380 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4381 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4382 : Intrinsic::x86_avx512_psll_q_512;
4383 else if (
Size ==
'w')
4384 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4385 : Intrinsic::x86_avx512_psll_w_512;
4391 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4392 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4393 bool IsVariable = Name[16] ==
'v';
4394 char Size = Name[16] ==
'.' ? Name[17]
4395 : Name[17] ==
'.' ? Name[18]
4396 : Name[18] ==
'.' ? Name[19]
4400 if (IsVariable && Name[17] !=
'.') {
4401 if (
Size ==
'd' && Name[17] ==
'2')
4402 IID = Intrinsic::x86_avx2_psrlv_q;
4403 else if (
Size ==
'd' && Name[17] ==
'4')
4404 IID = Intrinsic::x86_avx2_psrlv_q_256;
4405 else if (
Size ==
's' && Name[17] ==
'4')
4406 IID = Intrinsic::x86_avx2_psrlv_d;
4407 else if (
Size ==
's' && Name[17] ==
'8')
4408 IID = Intrinsic::x86_avx2_psrlv_d_256;
4409 else if (
Size ==
'h' && Name[17] ==
'8')
4410 IID = Intrinsic::x86_avx512_psrlv_w_128;
4411 else if (
Size ==
'h' && Name[17] ==
'1')
4412 IID = Intrinsic::x86_avx512_psrlv_w_256;
4413 else if (Name[17] ==
'3' && Name[18] ==
'2')
4414 IID = Intrinsic::x86_avx512_psrlv_w_512;
4417 }
else if (Name.ends_with(
".128")) {
4419 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4420 : Intrinsic::x86_sse2_psrl_d;
4421 else if (
Size ==
'q')
4422 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4423 : Intrinsic::x86_sse2_psrl_q;
4424 else if (
Size ==
'w')
4425 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4426 : Intrinsic::x86_sse2_psrl_w;
4429 }
else if (Name.ends_with(
".256")) {
4431 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4432 : Intrinsic::x86_avx2_psrl_d;
4433 else if (
Size ==
'q')
4434 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4435 : Intrinsic::x86_avx2_psrl_q;
4436 else if (
Size ==
'w')
4437 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4438 : Intrinsic::x86_avx2_psrl_w;
4443 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4444 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4445 : Intrinsic::x86_avx512_psrl_d_512;
4446 else if (
Size ==
'q')
4447 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4448 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4449 : Intrinsic::x86_avx512_psrl_q_512;
4450 else if (
Size ==
'w')
4451 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4452 : Intrinsic::x86_avx512_psrl_w_512;
4458 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4459 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4460 bool IsVariable = Name[16] ==
'v';
4461 char Size = Name[16] ==
'.' ? Name[17]
4462 : Name[17] ==
'.' ? Name[18]
4463 : Name[18] ==
'.' ? Name[19]
4467 if (IsVariable && Name[17] !=
'.') {
4468 if (
Size ==
's' && Name[17] ==
'4')
4469 IID = Intrinsic::x86_avx2_psrav_d;
4470 else if (
Size ==
's' && Name[17] ==
'8')
4471 IID = Intrinsic::x86_avx2_psrav_d_256;
4472 else if (
Size ==
'h' && Name[17] ==
'8')
4473 IID = Intrinsic::x86_avx512_psrav_w_128;
4474 else if (
Size ==
'h' && Name[17] ==
'1')
4475 IID = Intrinsic::x86_avx512_psrav_w_256;
4476 else if (Name[17] ==
'3' && Name[18] ==
'2')
4477 IID = Intrinsic::x86_avx512_psrav_w_512;
4480 }
else if (Name.ends_with(
".128")) {
4482 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4483 : Intrinsic::x86_sse2_psra_d;
4484 else if (
Size ==
'q')
4485 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4486 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4487 : Intrinsic::x86_avx512_psra_q_128;
4488 else if (
Size ==
'w')
4489 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4490 : Intrinsic::x86_sse2_psra_w;
4493 }
else if (Name.ends_with(
".256")) {
4495 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4496 : Intrinsic::x86_avx2_psra_d;
4497 else if (
Size ==
'q')
4498 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4499 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4500 : Intrinsic::x86_avx512_psra_q_256;
4501 else if (
Size ==
'w')
4502 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4503 : Intrinsic::x86_avx2_psra_w;
4508 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4509 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4510 : Intrinsic::x86_avx512_psra_d_512;
4511 else if (
Size ==
'q')
4512 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4513 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4514 : Intrinsic::x86_avx512_psra_q_512;
4515 else if (
Size ==
'w')
4516 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4517 : Intrinsic::x86_avx512_psra_w_512;
4523 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4525 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4527 }
else if (Name.ends_with(
".movntdqa")) {
4531 LoadInst *LI = Builder.CreateAlignedLoad(
4536 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4537 Name.starts_with(
"fma.vfmsub.") ||
4538 Name.starts_with(
"fma.vfnmadd.") ||
4539 Name.starts_with(
"fma.vfnmsub.")) {
4540 bool NegMul = Name[6] ==
'n';
4541 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4542 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4553 if (NegMul && !IsScalar)
4554 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4555 if (NegMul && IsScalar)
4556 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4558 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4560 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4564 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4572 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4576 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4577 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4578 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4579 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4580 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4581 bool IsMask3 = Name[11] ==
'3';
4582 bool IsMaskZ = Name[11] ==
'z';
4584 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4585 bool NegMul = Name[2] ==
'n';
4586 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4592 if (NegMul && (IsMask3 || IsMaskZ))
4593 A = Builder.CreateFNeg(
A);
4594 if (NegMul && !(IsMask3 || IsMaskZ))
4595 B = Builder.CreateFNeg(
B);
4597 C = Builder.CreateFNeg(
C);
4599 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4600 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4601 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4608 if (Name.back() ==
'd')
4609 IID = Intrinsic::x86_avx512_vfmadd_f64;
4611 IID = Intrinsic::x86_avx512_vfmadd_f32;
4612 Rep = Builder.CreateIntrinsic(IID,
Ops);
4614 Rep = Builder.CreateFMA(
A,
B,
C);
4623 if (NegAcc && IsMask3)
4628 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4630 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4631 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4632 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4633 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4634 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4635 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4636 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4637 bool IsMask3 = Name[11] ==
'3';
4638 bool IsMaskZ = Name[11] ==
'z';
4640 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4641 bool NegMul = Name[2] ==
'n';
4642 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4648 if (NegMul && (IsMask3 || IsMaskZ))
4649 A = Builder.CreateFNeg(
A);
4650 if (NegMul && !(IsMask3 || IsMaskZ))
4651 B = Builder.CreateFNeg(
B);
4653 C = Builder.CreateFNeg(
C);
4660 if (Name[Name.size() - 5] ==
's')
4661 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4663 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4667 Rep = Builder.CreateFMA(
A,
B,
C);
4675 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4679 if (VecWidth == 128 && EltWidth == 32)
4680 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4681 else if (VecWidth == 256 && EltWidth == 32)
4682 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4683 else if (VecWidth == 128 && EltWidth == 64)
4684 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4685 else if (VecWidth == 256 && EltWidth == 64)
4686 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4692 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4693 Rep = Builder.CreateIntrinsic(IID,
Ops);
4694 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4695 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4696 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4697 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4698 bool IsMask3 = Name[11] ==
'3';
4699 bool IsMaskZ = Name[11] ==
'z';
4701 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4702 bool IsSubAdd = Name[3] ==
's';
4706 if (Name[Name.size() - 5] ==
's')
4707 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4709 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4714 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4716 Rep = Builder.CreateIntrinsic(IID,
Ops);
4725 Value *Odd = Builder.CreateCall(FMA,
Ops);
4726 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4727 Value *Even = Builder.CreateCall(FMA,
Ops);
4733 for (
int i = 0; i != NumElts; ++i)
4734 Idxs[i] = i + (i % 2) * NumElts;
4736 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4744 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4745 Name.starts_with(
"avx512.maskz.pternlog.")) {
4746 bool ZeroMask = Name[11] ==
'z';
4750 if (VecWidth == 128 && EltWidth == 32)
4751 IID = Intrinsic::x86_avx512_pternlog_d_128;
4752 else if (VecWidth == 256 && EltWidth == 32)
4753 IID = Intrinsic::x86_avx512_pternlog_d_256;
4754 else if (VecWidth == 512 && EltWidth == 32)
4755 IID = Intrinsic::x86_avx512_pternlog_d_512;
4756 else if (VecWidth == 128 && EltWidth == 64)
4757 IID = Intrinsic::x86_avx512_pternlog_q_128;
4758 else if (VecWidth == 256 && EltWidth == 64)
4759 IID = Intrinsic::x86_avx512_pternlog_q_256;
4760 else if (VecWidth == 512 && EltWidth == 64)
4761 IID = Intrinsic::x86_avx512_pternlog_q_512;
4767 Rep = Builder.CreateIntrinsic(IID, Args);
4771 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4772 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4773 bool ZeroMask = Name[11] ==
'z';
4774 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4777 if (VecWidth == 128 && !
High)
4778 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4779 else if (VecWidth == 256 && !
High)
4780 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4781 else if (VecWidth == 512 && !
High)
4782 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4783 else if (VecWidth == 128 &&
High)
4784 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4785 else if (VecWidth == 256 &&
High)
4786 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4787 else if (VecWidth == 512 &&
High)
4788 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4794 Rep = Builder.CreateIntrinsic(IID, Args);
4798 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4799 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4800 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4801 bool ZeroMask = Name[11] ==
'z';
4802 bool IndexForm = Name[17] ==
'i';
4804 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4805 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4806 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4807 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4808 bool ZeroMask = Name[11] ==
'z';
4809 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4812 if (VecWidth == 128 && !IsSaturating)
4813 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4814 else if (VecWidth == 256 && !IsSaturating)
4815 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4816 else if (VecWidth == 512 && !IsSaturating)
4817 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4818 else if (VecWidth == 128 && IsSaturating)
4819 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4820 else if (VecWidth == 256 && IsSaturating)
4821 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4822 else if (VecWidth == 512 && IsSaturating)
4823 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4833 if (Args[1]->
getType()->isVectorTy() &&
4836 ->isIntegerTy(32) &&
4837 Args[2]->
getType()->isVectorTy() &&
4840 ->isIntegerTy(32)) {
4841 Type *NewArgType =
nullptr;
4842 if (VecWidth == 128)
4844 else if (VecWidth == 256)
4846 else if (VecWidth == 512)
4852 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4853 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4856 Rep = Builder.CreateIntrinsic(IID, Args);
4860 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4861 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4862 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4863 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4864 bool ZeroMask = Name[11] ==
'z';
4865 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4868 if (VecWidth == 128 && !IsSaturating)
4869 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4870 else if (VecWidth == 256 && !IsSaturating)
4871 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4872 else if (VecWidth == 512 && !IsSaturating)
4873 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4874 else if (VecWidth == 128 && IsSaturating)
4875 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4876 else if (VecWidth == 256 && IsSaturating)
4877 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4878 else if (VecWidth == 512 && IsSaturating)
4879 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4889 if (Args[1]->
getType()->isVectorTy() &&
4892 ->isIntegerTy(32) &&
4893 Args[2]->
getType()->isVectorTy() &&
4896 ->isIntegerTy(32)) {
4897 Type *NewArgType =
nullptr;
4898 if (VecWidth == 128)
4900 else if (VecWidth == 256)
4902 else if (VecWidth == 512)
4908 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4909 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4912 Rep = Builder.CreateIntrinsic(IID, Args);
4916 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4917 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4918 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4920 if (Name[0] ==
'a' && Name.back() ==
'2')
4921 IID = Intrinsic::x86_addcarry_32;
4922 else if (Name[0] ==
'a' && Name.back() ==
'4')
4923 IID = Intrinsic::x86_addcarry_64;
4924 else if (Name[0] ==
's' && Name.back() ==
'2')
4925 IID = Intrinsic::x86_subborrow_32;
4926 else if (Name[0] ==
's' && Name.back() ==
'4')
4927 IID = Intrinsic::x86_subborrow_64;
4934 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4937 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4940 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4944 }
else if (Name.starts_with(
"avx512.mask.") &&
4947 }
else if (Name.starts_with(
"bmi.pdep.")) {
4949 }
else if (Name.starts_with(
"bmi.pext.")) {
4959 if (Name.starts_with(
"neon.bfcvt")) {
4960 if (Name.starts_with(
"neon.bfcvtn2")) {
4962 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4964 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4965 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4968 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4969 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4971 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4975 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4976 return Builder.CreateShuffleVector(
4979 return Builder.CreateFPTrunc(CI->
getOperand(0),
4982 }
else if (Name.starts_with(
"sve.fcvt")) {
4985 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4986 .
Case(
"sve.fcvtnt.bf16f32",
4987 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4999 if (Args[1]->
getType() != BadPredTy)
5002 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5003 BadPredTy, Args[1]);
5004 Args[1] = Builder.CreateIntrinsic(
5005 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5007 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5011 if (Name ==
"neon.vcvtfp2hf")
5012 return Builder.CreateBitCast(
5013 Builder.CreateFPTrunc(
5017 if (Name ==
"neon.vcvthf2fp")
5018 return Builder.CreateFPExt(
5019 Builder.CreateBitCast(
5029 if (Name ==
"mve.vctp64.old") {
5032 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5035 Value *C1 = Builder.CreateIntrinsic(
5036 Intrinsic::arm_mve_pred_v2i,
5038 return Builder.CreateIntrinsic(
5039 Intrinsic::arm_mve_pred_i2v,
5041 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5042 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5043 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5044 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5046 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5047 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5048 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5049 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5051 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5052 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5053 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5054 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5055 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5056 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5057 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5058 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5059 std::vector<Type *> Tys;
5063 case Intrinsic::arm_mve_mull_int_predicated:
5064 case Intrinsic::arm_mve_vqdmull_predicated:
5065 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5068 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5069 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5070 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5074 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5078 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5082 case Intrinsic::arm_cde_vcx1q_predicated:
5083 case Intrinsic::arm_cde_vcx1qa_predicated:
5084 case Intrinsic::arm_cde_vcx2q_predicated:
5085 case Intrinsic::arm_cde_vcx2qa_predicated:
5086 case Intrinsic::arm_cde_vcx3q_predicated:
5087 case Intrinsic::arm_cde_vcx3qa_predicated:
5094 std::vector<Value *>
Ops;
5096 Type *Ty =
Op->getType();
5097 if (Ty->getScalarSizeInBits() == 1) {
5098 Value *C1 = Builder.CreateIntrinsic(
5099 Intrinsic::arm_mve_pred_v2i,
5101 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5106 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5121 auto UpgradeLegacyWMMAIUIntrinsicCall =
5126 Args.push_back(Builder.getFalse());
5130 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5137 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5141 NewCall->copyMetadata(*CI);
5145 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5146 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5147 "intrinsic should have 7 arguments");
5150 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5152 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5153 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5154 "intrinsic should have 8 arguments");
5159 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5162 switch (
F->getIntrinsicID()) {
5165 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5166 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5167 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5168 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5169 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5170 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5185 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5188 F->getParent(),
F->getIntrinsicID(), Overloads);
5193 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5197 NewCall->copyMetadata(*CI);
5198 NewCall->takeName(CI);
5203 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5209 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5210 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5235 if (NumOperands < 3)
5248 bool IsVolatile =
false;
5252 if (NumOperands > 3)
5257 if (NumOperands > 5) {
5259 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5273 if (VT->getElementType()->isIntegerTy(16)) {
5276 Val = Builder.CreateBitCast(Val, AsBF16);
5284 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5286 unsigned AddrSpace = PtrTy->getAddressSpace();
5289 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5291 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5296 MDNode *RangeNotPrivate =
5299 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5305 return Builder.CreateBitCast(RMW, RetTy);
5326 return MAV->getMetadata();
5335 if (Name ==
"label") {
5337 }
else if (Name ==
"assign") {
5344 }
else if (Name ==
"declare") {
5348 }
else if (Name ==
"addr") {
5358 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5359 }
else if (Name ==
"value") {
5362 unsigned ExprOp = 2;
5377 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5385 int64_t OffsetVal =
Offset->getSExtValue();
5386 return Builder.CreateIntrinsic(OffsetVal >= 0
5387 ? Intrinsic::vector_splice_left
5388 : Intrinsic::vector_splice_right,
5390 {CI->getArgOperand(0), CI->getArgOperand(1),
5391 Builder.getInt32(std::abs(OffsetVal))});
5396 if (Name.starts_with(
"to.fp16")) {
5398 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5399 return Builder.CreateBitCast(Cast, CI->
getType());
5402 if (Name.starts_with(
"from.fp16")) {
5404 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5405 return Builder.CreateFPExt(Cast, CI->
getType());
5464 else if (Opcode == Instruction::ICmp)
5467 else if (Opcode == Instruction::FCmp)
5470 else if (Opcode == Instruction::Select)
5475 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5487 if (Defaults.empty())
5490 unsigned OldArgCount = CI->
arg_size();
5491 unsigned NewArgCount = NewFn->
arg_size();
5495 if (OldArgCount >= NewArgCount)
5503 if (OldArgCount < FirstDefault)
5508 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5509 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5510 "missing argument outside the default range");
5511 Type *ParamTy = NewFT->getParamType(Idx);
5516 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5522 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5554 if (!Name.consume_front(
"llvm."))
5557 bool IsX86 = Name.consume_front(
"x86.");
5558 bool IsNVVM = Name.consume_front(
"nvvm.");
5559 bool IsAArch64 = Name.consume_front(
"aarch64.");
5560 bool IsARM = Name.consume_front(
"arm.");
5561 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5562 bool IsDbg = Name.consume_front(
"dbg.");
5564 (Name.consume_front(
"experimental.vector.splice") ||
5565 Name.consume_front(
"vector.splice")) &&
5566 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5567 Value *Rep =
nullptr;
5569 if (!IsX86 && Name ==
"stackprotectorcheck") {
5571 }
else if (IsNVVM) {
5575 }
else if (IsAArch64) {
5579 }
else if (IsAMDGCN) {
5583 }
else if (IsOldSplice) {
5585 }
else if (Name.consume_front(
"convert.")) {
5587 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5602 const auto &DefaultCase = [&]() ->
void {
5610 "Unknown function for CallBase upgrade and isn't just a name change");
5618 "Return type must have changed");
5619 assert(OldST->getNumElements() ==
5621 "Must have same number of elements");
5624 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5627 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5628 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5629 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5653 case Intrinsic::arm_neon_vst1:
5654 case Intrinsic::arm_neon_vst2:
5655 case Intrinsic::arm_neon_vst3:
5656 case Intrinsic::arm_neon_vst4:
5657 case Intrinsic::arm_neon_vst2lane:
5658 case Intrinsic::arm_neon_vst3lane:
5659 case Intrinsic::arm_neon_vst4lane: {
5661 NewCall = Builder.CreateCall(NewFn, Args);
5664 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5665 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5666 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5671 NewCall = Builder.CreateCall(NewFn, Args);
5674 case Intrinsic::aarch64_sve_ld3_sret:
5675 case Intrinsic::aarch64_sve_ld4_sret:
5676 case Intrinsic::aarch64_sve_ld2_sret: {
5684 Name = Name.substr(5);
5691 unsigned MinElts = RetTy->getMinNumElements() /
N;
5693 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5695 for (
unsigned I = 0;
I <
N;
I++) {
5696 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5697 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5703 case Intrinsic::coro_end_async:
5704 case Intrinsic::coro_end: {
5706 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5708 NewCall = Builder.CreateCall(NewFn, Args);
5713 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5714 Value *InRamp = Builder.CreateCall(IsInRamp);
5724 case Intrinsic::vector_extract: {
5726 Name = Name.substr(5);
5727 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5732 unsigned MinElts = RetTy->getMinNumElements();
5735 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5739 case Intrinsic::vector_insert: {
5741 Name = Name.substr(5);
5742 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5746 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5751 NewCall = Builder.CreateCall(
5755 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5761 assert(
N > 1 &&
"Create is expected to be between 2-4");
5764 unsigned MinElts = RetTy->getMinNumElements() /
N;
5765 for (
unsigned I = 0;
I <
N;
I++) {
5767 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5774 case Intrinsic::arm_neon_bfdot:
5775 case Intrinsic::arm_neon_bfmmla:
5776 case Intrinsic::arm_neon_bfmlalb:
5777 case Intrinsic::arm_neon_bfmlalt:
5778 case Intrinsic::aarch64_neon_bfdot:
5779 case Intrinsic::aarch64_neon_bfmmla:
5780 case Intrinsic::aarch64_neon_bfmlalb:
5781 case Intrinsic::aarch64_neon_bfmlalt: {
5784 "Mismatch between function args and call args");
5785 size_t OperandWidth =
5787 assert((OperandWidth == 64 || OperandWidth == 128) &&
5788 "Unexpected operand width");
5790 auto Iter = CI->
args().begin();
5791 Args.push_back(*Iter++);
5792 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5793 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5794 NewCall = Builder.CreateCall(NewFn, Args);
5798 case Intrinsic::bitreverse:
5799 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5802 case Intrinsic::ctlz:
5803 case Intrinsic::cttz: {
5810 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5814 case Intrinsic::objectsize: {
5815 Value *NullIsUnknownSize =
5819 NewCall = Builder.CreateCall(
5824 case Intrinsic::ctpop:
5825 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5827 case Intrinsic::dbg_value: {
5829 Name = Name.substr(5);
5831 if (Name.starts_with(
"dbg.addr")) {
5845 if (
Offset->isNullValue()) {
5846 NewCall = Builder.CreateCall(
5855 case Intrinsic::ptr_annotation:
5863 NewCall = Builder.CreateCall(
5872 case Intrinsic::var_annotation:
5879 NewCall = Builder.CreateCall(
5888 case Intrinsic::riscv_aes32dsi:
5889 case Intrinsic::riscv_aes32dsmi:
5890 case Intrinsic::riscv_aes32esi:
5891 case Intrinsic::riscv_aes32esmi:
5892 case Intrinsic::riscv_sm4ks:
5893 case Intrinsic::riscv_sm4ed: {
5903 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5904 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5910 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5911 Value *Res = NewCall;
5913 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5919 case Intrinsic::nvvm_mapa_shared_cluster: {
5923 Value *Res = NewCall;
5924 Res = Builder.CreateAddrSpaceCast(
5931 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5932 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5935 Args[0] = Builder.CreateAddrSpaceCast(
5938 NewCall = Builder.CreateCall(NewFn, Args);
5944 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5945 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5946 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5947 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5948 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5949 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5950 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5951 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5958 Args[0] = Builder.CreateAddrSpaceCast(
5967 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5969 NewCall = Builder.CreateCall(NewFn, Args);
5975 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
5976 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
5977 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
5978 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
5979 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
5980 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
5981 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
5982 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
5984 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
5988 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
5989 NewCall = Builder.CreateCall(NewFn, Args);
5992 case Intrinsic::nvvm_tcgen05_mma_shared:
5993 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5994 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5995 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
5996 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
5997 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
5998 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
5999 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
6000 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
6001 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
6002 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
6003 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
6004 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
6005 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
6006 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
6007 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
6008 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
6009 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
6010 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
6011 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
6012 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
6013 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
6014 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
6015 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
6016 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
6017 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
6018 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
6019 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
6020 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
6021 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
6022 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
6023 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
6024 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
6025 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
6026 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
6027 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
6028 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
6029 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
6030 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
6032 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
6033 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
6035 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
6036 case Intrinsic::nvvm_tcgen05_mma_tensor:
6037 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
6038 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
6039 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
6040 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
6041 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
6042 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
6043 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
6044 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
6045 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
6046 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
6047 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
6048 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
6049 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
6050 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
6052 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
6053 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
6055 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
6057 Args.push_back(Builder.getInt32(0));
6058 NewCall = Builder.CreateCall(NewFn, Args);
6061 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6062 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6063 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6064 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6067 Builder.getFalse()});
6069 case Intrinsic::riscv_sha256sig0:
6070 case Intrinsic::riscv_sha256sig1:
6071 case Intrinsic::riscv_sha256sum0:
6072 case Intrinsic::riscv_sha256sum1:
6073 case Intrinsic::riscv_sm3p0:
6074 case Intrinsic::riscv_sm3p1: {
6081 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6083 NewCall = Builder.CreateCall(NewFn, Arg);
6085 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6092 case Intrinsic::x86_xop_vfrcz_ss:
6093 case Intrinsic::x86_xop_vfrcz_sd:
6094 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6097 case Intrinsic::x86_xop_vpermil2pd:
6098 case Intrinsic::x86_xop_vpermil2ps:
6099 case Intrinsic::x86_xop_vpermil2pd_256:
6100 case Intrinsic::x86_xop_vpermil2ps_256: {
6104 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6105 NewCall = Builder.CreateCall(NewFn, Args);
6109 case Intrinsic::x86_sse41_ptestc:
6110 case Intrinsic::x86_sse41_ptestz:
6111 case Intrinsic::x86_sse41_ptestnzc: {
6125 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6126 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6128 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6132 case Intrinsic::x86_rdtscp: {
6138 NewCall = Builder.CreateCall(NewFn);
6140 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6143 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6151 case Intrinsic::x86_sse41_insertps:
6152 case Intrinsic::x86_sse41_dppd:
6153 case Intrinsic::x86_sse41_dpps:
6154 case Intrinsic::x86_sse41_mpsadbw:
6155 case Intrinsic::x86_avx_dp_ps_256:
6156 case Intrinsic::x86_avx2_mpsadbw: {
6162 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6163 NewCall = Builder.CreateCall(NewFn, Args);
6167 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6168 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6169 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6170 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6171 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6172 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6178 NewCall = Builder.CreateCall(NewFn, Args);
6187 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6188 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6189 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6190 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6191 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6192 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6196 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6197 Args[1] = Builder.CreateBitCast(
6200 NewCall = Builder.CreateCall(NewFn, Args);
6201 Value *Res = Builder.CreateBitCast(
6209 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6210 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6211 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6215 Args[1] = Builder.CreateBitCast(
6217 Args[2] = Builder.CreateBitCast(
6220 NewCall = Builder.CreateCall(NewFn, Args);
6224 case Intrinsic::thread_pointer: {
6225 NewCall = Builder.CreateCall(NewFn, {});
6229 case Intrinsic::memcpy:
6230 case Intrinsic::memmove:
6231 case Intrinsic::memset: {
6247 NewCall = Builder.CreateCall(NewFn, Args);
6249 AttributeList NewAttrs = AttributeList::get(
6250 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6251 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6252 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6257 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6260 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6264 case Intrinsic::masked_load:
6265 case Intrinsic::masked_gather:
6266 case Intrinsic::masked_store:
6267 case Intrinsic::masked_scatter: {
6273 auto GetMaybeAlign = [](
Value *
Op) {
6275 uint64_t Val = CI->getZExtValue();
6283 auto GetAlign = [&](
Value *
Op) {
6292 case Intrinsic::masked_load:
6293 NewCall = Builder.CreateMaskedLoad(
6297 case Intrinsic::masked_gather:
6298 NewCall = Builder.CreateMaskedGather(
6304 case Intrinsic::masked_store:
6305 NewCall = Builder.CreateMaskedStore(
6309 case Intrinsic::masked_scatter:
6310 NewCall = Builder.CreateMaskedScatter(
6312 DL.getValueOrABITypeAlignment(
6326 case Intrinsic::lifetime_start:
6327 case Intrinsic::lifetime_end: {
6339 NewCall = Builder.CreateLifetimeStart(Ptr);
6341 NewCall = Builder.CreateLifetimeEnd(Ptr);
6350 case Intrinsic::x86_avx512_vpdpbusd_128:
6351 case Intrinsic::x86_avx512_vpdpbusd_256:
6352 case Intrinsic::x86_avx512_vpdpbusd_512:
6353 case Intrinsic::x86_avx512_vpdpbusds_128:
6354 case Intrinsic::x86_avx512_vpdpbusds_256:
6355 case Intrinsic::x86_avx512_vpdpbusds_512:
6356 case Intrinsic::x86_avx2_vpdpbssd_128:
6357 case Intrinsic::x86_avx2_vpdpbssd_256:
6358 case Intrinsic::x86_avx10_vpdpbssd_512:
6359 case Intrinsic::x86_avx2_vpdpbssds_128:
6360 case Intrinsic::x86_avx2_vpdpbssds_256:
6361 case Intrinsic::x86_avx10_vpdpbssds_512:
6362 case Intrinsic::x86_avx2_vpdpbsud_128:
6363 case Intrinsic::x86_avx2_vpdpbsud_256:
6364 case Intrinsic::x86_avx10_vpdpbsud_512:
6365 case Intrinsic::x86_avx2_vpdpbsuds_128:
6366 case Intrinsic::x86_avx2_vpdpbsuds_256:
6367 case Intrinsic::x86_avx10_vpdpbsuds_512:
6368 case Intrinsic::x86_avx2_vpdpbuud_128:
6369 case Intrinsic::x86_avx2_vpdpbuud_256:
6370 case Intrinsic::x86_avx10_vpdpbuud_512:
6371 case Intrinsic::x86_avx2_vpdpbuuds_128:
6372 case Intrinsic::x86_avx2_vpdpbuuds_256:
6373 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6378 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6379 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6381 NewCall = Builder.CreateCall(NewFn, Args);
6384 case Intrinsic::x86_avx512_vpdpwssd_128:
6385 case Intrinsic::x86_avx512_vpdpwssd_256:
6386 case Intrinsic::x86_avx512_vpdpwssd_512:
6387 case Intrinsic::x86_avx512_vpdpwssds_128:
6388 case Intrinsic::x86_avx512_vpdpwssds_256:
6389 case Intrinsic::x86_avx512_vpdpwssds_512:
6390 case Intrinsic::x86_avx2_vpdpwsud_128:
6391 case Intrinsic::x86_avx2_vpdpwsud_256:
6392 case Intrinsic::x86_avx10_vpdpwsud_512:
6393 case Intrinsic::x86_avx2_vpdpwsuds_128:
6394 case Intrinsic::x86_avx2_vpdpwsuds_256:
6395 case Intrinsic::x86_avx10_vpdpwsuds_512:
6396 case Intrinsic::x86_avx2_vpdpwusd_128:
6397 case Intrinsic::x86_avx2_vpdpwusd_256:
6398 case Intrinsic::x86_avx10_vpdpwusd_512:
6399 case Intrinsic::x86_avx2_vpdpwusds_128:
6400 case Intrinsic::x86_avx2_vpdpwusds_256:
6401 case Intrinsic::x86_avx10_vpdpwusds_512:
6402 case Intrinsic::x86_avx2_vpdpwuud_128:
6403 case Intrinsic::x86_avx2_vpdpwuud_256:
6404 case Intrinsic::x86_avx10_vpdpwuud_512:
6405 case Intrinsic::x86_avx2_vpdpwuuds_128:
6406 case Intrinsic::x86_avx2_vpdpwuuds_256:
6407 case Intrinsic::x86_avx10_vpdpwuuds_512:
6412 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6413 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6415 NewCall = Builder.CreateCall(NewFn, Args);
6418 assert(NewCall &&
"Should have either set this variable or returned through "
6419 "the default case");
6426 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6440 F->eraseFromParent();
6446 if (NumOperands == 0)
6454 if (NumOperands == 3) {
6458 Metadata *Elts2[] = {ScalarType, ScalarType,
6472 if (
Opc != Instruction::BitCast)
6476 Type *SrcTy = V->getType();
6493 if (
Opc != Instruction::BitCast)
6496 Type *SrcTy =
C->getType();
6513 if (Flag.getNumOperands() < 3)
6514 return std::nullopt;
6516 return Name->getString();
6517 return std::nullopt;
6531 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6532 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6533 if (auto Name = getModuleFlagNameSafely(*Flag))
6534 return *Name ==
"Debug Info Version";
6537 if (OpIt != ModFlags->op_end()) {
6538 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6545 bool BrokenDebugInfo =
false;
6548 if (!BrokenDebugInfo)
6554 M.getContext().diagnose(Diag);
6561 M.getContext().diagnose(DiagVersion);
6571 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6574 if (
F->hasFnAttribute(Attr)) {
6577 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6579 auto [Part, Rest] = S.
split(
',');
6585 const unsigned Dim = DimC -
'x';
6586 assert(Dim < 3 &&
"Unexpected dim char");
6596 F->addFnAttr(Attr, NewAttr);
6600 return S ==
"x" || S ==
"y" || S ==
"z";
6605 if (K ==
"kernel") {
6617 const unsigned Idx = (AlignIdxValuePair >> 16);
6618 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6623 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6628 if (K ==
"minctasm") {
6633 if (K ==
"maxnreg") {
6638 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6642 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6646 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6650 if (K ==
"grid_constant") {
6665 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6672 if (!SeenNodes.
insert(MD).second)
6679 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6686 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6688 const MDOperand &V = MD->getOperand(j + 1);
6691 NewOperands.
append({K, V});
6694 if (NewOperands.
size() > 1)
6707 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6708 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6709 if (ModRetainReleaseMarker) {
6715 ID->getString().split(ValueComp,
"#");
6716 if (ValueComp.
size() == 2) {
6717 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6721 M.eraseNamedMetadata(ModRetainReleaseMarker);
6732 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6758 bool InvalidCast =
false;
6760 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6773 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6775 Args.push_back(Arg);
6782 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6787 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6800 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6808 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6809 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6810 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6811 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6812 {
"objc_autoreleaseReturnValue",
6813 llvm::Intrinsic::objc_autoreleaseReturnValue},
6814 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6815 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6816 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6817 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6818 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6819 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6820 {
"objc_release", llvm::Intrinsic::objc_release},
6821 {
"objc_retain", llvm::Intrinsic::objc_retain},
6822 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6823 {
"objc_retainAutoreleaseReturnValue",
6824 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6825 {
"objc_retainAutoreleasedReturnValue",
6826 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6827 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6828 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6829 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6830 {
"objc_unsafeClaimAutoreleasedReturnValue",
6831 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6832 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6833 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6834 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6835 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6836 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6837 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6838 {
"objc_arc_annotation_topdown_bbstart",
6839 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6840 {
"objc_arc_annotation_topdown_bbend",
6841 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6842 {
"objc_arc_annotation_bottomup_bbstart",
6843 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6844 {
"objc_arc_annotation_bottomup_bbend",
6845 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6847 for (
auto &
I : RuntimeFuncs)
6848 UpgradeToIntrinsic(
I.first,
I.second);
6872 std::optional<bool> UseAddressDisc;
6875 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6876 for (
const MDNode *Flag : ModFlags->operands()) {
6878 if (Name && (*Name ==
"ptrauth-init-fini" ||
6879 *Name ==
"ptrauth-init-fini-address-discrimination"))
6884 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6885 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6886 constexpr unsigned ExpectedAddressMarker = 1;
6889 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6892 bool HasAddressDisc;
6893 if (!CPA->hasAddressDiscriminator())
6894 HasAddressDisc =
false;
6895 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6896 HasAddressDisc =
true;
6900 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6903 UseAddressDisc = HasAddressDisc;
6904 return CPA->getPointer();
6908 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6911 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6913 if (!GV || !GV->hasInitializer())
6917 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6920 std::vector<Constant *> NewStructors;
6921 NewStructors.reserve(OldStructorsArray->getNumOperands());
6923 for (
Use &U : OldStructorsArray->operands()) {
6932 Func = UpgradeSinglePointer(Func);
6936 NewStructors.push_back(
6945 if (GlobalArraysToUpgrade.
empty())
6947 assert(UseAddressDisc.has_value());
6949 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6950 GV->setInitializer(NewInit);
6953 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6963 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6967 bool HasObjCFlag =
false, HasClassProperties =
false;
6968 bool HasSwiftVersionFlag =
false;
6969 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6976 if (
Op->getNumOperands() != 3)
6990 if (ID->getString() ==
"Objective-C Image Info Version")
6992 if (ID->getString() ==
"Objective-C Class Properties")
6993 HasClassProperties =
true;
6995 if (ID->getString() ==
"PIC Level") {
6996 if (
auto *Behavior =
6998 uint64_t V = Behavior->getLimitedValue();
7004 if (ID->getString() ==
"PIE Level")
7005 if (
auto *Behavior =
7012 if (ID->getString() ==
"branch-target-enforcement" ||
7013 ID->getString().starts_with(
"sign-return-address")) {
7014 if (
auto *Behavior =
7020 Op->getOperand(1),
Op->getOperand(2)};
7030 if (ID->getString() ==
"Objective-C Image Info Section") {
7033 Value->getString().split(ValueComp,
" ");
7034 if (ValueComp.
size() != 1) {
7035 std::string NewValue;
7036 for (
auto &S : ValueComp)
7037 NewValue += S.str();
7048 if (ID->getString() ==
"Objective-C Garbage Collection") {
7051 assert(Md->getValue() &&
"Expected non-empty metadata");
7052 auto Type = Md->getValue()->getType();
7055 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7056 if ((Val & 0xff) != Val) {
7057 HasSwiftVersionFlag =
true;
7058 SwiftABIVersion = (Val & 0xff00) >> 8;
7059 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7060 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7071 if (ID->getString() ==
"amdgpu_code_object_version") {
7074 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7083 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7112 if (HasObjCFlag && !HasClassProperties) {
7118 if (HasSwiftVersionFlag) {
7122 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7124 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7132 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7136 auto MatchesVersion = [](
const MDNode *
Op) {
7137 return Op->getNumOperands() >= 3 &&
7151 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7152 assert(
Op->getNumOperands() >= 2 &&
7153 "Expected at least 2 operands - name and linkage type");
7165 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7176 auto TrimSpaces = [](
StringRef Section) -> std::string {
7178 Section.split(Components,
',');
7183 for (
auto Component : Components)
7184 OS <<
',' << Component.trim();
7189 for (
auto &GV : M.globals()) {
7190 if (!GV.hasSection())
7195 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7200 GV.setSection(TrimSpaces(Section));
7216struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7217 StrictFPUpgradeVisitor() =
default;
7220 if (!
Call.isStrictFP())
7226 Call.removeFnAttr(Attribute::StrictFP);
7227 Call.addFnAttr(Attribute::NoBuiltin);
7232struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7233 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7234 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7236 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7251 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7252 StrictFPUpgradeVisitor SFPV;
7257 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7258 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7259 for (
auto &Arg :
F.args())
7261 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7263 bool AddingAttrs =
false, RemovingAttrs =
false;
7264 AttrBuilder AttrsToAdd(
F.getContext());
7269 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7270 A.isValid() &&
A.isStringAttribute()) {
7271 F.setSection(
A.getValueAsString());
7273 RemovingAttrs =
true;
7277 A.isValid() &&
A.isStringAttribute()) {
7280 AddingAttrs = RemovingAttrs =
true;
7283 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7284 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7286 RemovingAttrs =
true;
7287 if (
A.getValueAsString() ==
"true") {
7288 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7297 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7300 if (
A.getValueAsBool()) {
7301 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7307 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7308 RemovingAttrs =
true;
7315 bool HandleDenormalMode =
false;
7317 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7320 DenormalFPMath = ParsedMode;
7322 AddingAttrs = RemovingAttrs =
true;
7323 HandleDenormalMode =
true;
7327 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7331 DenormalFPMathF32 = ParsedMode;
7333 AddingAttrs = RemovingAttrs =
true;
7334 HandleDenormalMode =
true;
7338 if (HandleDenormalMode)
7339 AttrsToAdd.addDenormalFPEnvAttr(
7343 F.removeFnAttrs(AttrsToRemove);
7346 F.addFnAttrs(AttrsToAdd);
7352 if (!
F.hasFnAttribute(FnAttrName))
7353 F.addFnAttr(FnAttrName,
Value);
7360 if (!
F.hasFnAttribute(FnAttrName)) {
7362 F.addFnAttr(FnAttrName);
7364 auto A =
F.getFnAttribute(FnAttrName);
7365 if (
"false" ==
A.getValueAsString())
7366 F.removeFnAttr(FnAttrName);
7367 else if (
"true" ==
A.getValueAsString()) {
7368 F.removeFnAttr(FnAttrName);
7369 F.addFnAttr(FnAttrName);
7375 Triple T(M.getTargetTriple());
7376 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7379 uint64_t BTEValue = 0;
7380 uint64_t BPPLRValue = 0;
7381 uint64_t GCSValue = 0;
7382 uint64_t SRAValue = 0;
7383 uint64_t SRAALLValue = 0;
7384 uint64_t SRABKeyValue = 0;
7386 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7390 if (
Op->getNumOperands() != 3)
7399 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7400 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7401 : IDStr ==
"guarded-control-stack" ? &GCSValue
7402 : IDStr ==
"sign-return-address" ? &SRAValue
7403 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7404 : IDStr ==
"sign-return-address-with-bkey"
7410 *ValPtr = CI->getZExtValue();
7416 bool BTE = BTEValue == 1;
7417 bool BPPLR = BPPLRValue == 1;
7418 bool GCS = GCSValue == 1;
7419 bool SRA = SRAValue == 1;
7422 if (SRA && SRAALLValue == 1)
7423 SignTypeValue =
"all";
7426 if (SRA && SRABKeyValue == 1)
7427 SignKeyValue =
"b_key";
7429 for (
Function &
F : M.getFunctionList()) {
7430 if (
F.isDeclaration())
7437 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7438 A.isValid() &&
"none" ==
A.getValueAsString()) {
7439 F.removeFnAttr(
"sign-return-address");
7440 F.removeFnAttr(
"sign-return-address-key");
7456 if (SRAALLValue == 1)
7458 if (SRABKeyValue == 1)
7485 if (
T->getNumOperands() < 1)
7490 if (S->getString().starts_with(
"llvm.vectorizer."))
7496 StringRef OldPrefix =
"llvm.vectorizer.";
7499 if (OldTag ==
"llvm.vectorizer.unroll")
7511 if (
T->getNumOperands() < 1)
7523 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7536 Ops.reserve(
T->getNumOperands());
7537 Ops.push_back(NewTag);
7538 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7539 Ops.push_back(
T->getOperand(
I));
7556 if (
T->isDistinct()) {
7557 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7569 Ops.reserve(
T->getNumOperands());
7580 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7581 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7582 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7585 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7587 auto I =
DL.find(
"-n64-");
7589 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7594 std::string Res =
DL.str();
7597 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7598 Res.append(Res.empty() ?
"G1" :
"-G1");
7606 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7607 Res.append(
"-ni:7:8:9");
7609 if (
DL.ends_with(
"ni:7"))
7611 if (
DL.ends_with(
"ni:7:8"))
7616 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7617 Res.append(
"-p7:160:256:256:32");
7618 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7619 Res.append(
"-p8:128:128:128:48");
7620 constexpr StringRef OldP8(
"-p8:128:128-");
7621 if (
DL.contains(OldP8))
7622 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7623 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7624 Res.append(
"-p9:192:256:256:32");
7628 if (!
DL.contains(
"m:e"))
7629 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7634 if (
T.isSystemZ() && !
DL.empty()) {
7636 if (!
DL.contains(
"-S64"))
7637 return "E-S64" +
DL.drop_front(1).str();
7641 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7644 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7645 if (!
DL.contains(AddrSpaces)) {
7647 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7648 if (R.match(Res, &
Groups))
7654 if (
T.isAArch64()) {
7656 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7657 Res.append(
"-Fn32");
7658 AddPtr32Ptr64AddrSpaces();
7662 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7666 std::string I64 =
"-i64:64";
7667 std::string I128 =
"-i128:128";
7669 size_t Pos = Res.find(I64);
7670 if (Pos !=
size_t(-1))
7671 Res.insert(Pos + I64.size(), I128);
7675 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7676 size_t Pos = Res.find(
"-S128");
7679 Res.insert(Pos,
"-f64:32:64");
7685 AddPtr32Ptr64AddrSpaces();
7693 if (!
T.isOSIAMCU()) {
7694 std::string I128 =
"-i128:128";
7697 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7698 if (R.match(Res, &
Groups))
7706 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7708 auto I =
Ref.find(
"-f80:32-");
7710 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7718 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7721 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7722 B.removeAttribute(
"no-frame-pointer-elim");
7724 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7726 if (FramePointer !=
"all")
7727 FramePointer =
"non-leaf";
7728 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7730 if (!FramePointer.
empty())
7731 B.addAttribute(
"frame-pointer", FramePointer);
7733 A =
B.getAttribute(
"null-pointer-is-valid");
7736 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7737 B.removeAttribute(
"null-pointer-is-valid");
7738 if (NullPointerIsValid)
7739 B.addAttribute(Attribute::NullPointerIsValid);
7742 A =
B.getAttribute(
"uniform-work-group-size");
7746 bool IsTrue = Val ==
"true";
7747 B.removeAttribute(
"uniform-work-group-size");
7749 B.addAttribute(
"uniform-work-group-size");
7760 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 bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
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 FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
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.