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)
1174#define NVVM_TMA_G2S_MODES(M) \
1175 M(tile_1d, "tile.1d") \
1176 M(tile_2d, "tile.2d") \
1177 M(tile_3d, "tile.3d") \
1178 M(tile_4d, "tile.4d") \
1179 M(tile_5d, "tile.5d") \
1180 M(tile_gather4_2d, "tile.gather4.2d") \
1181 M(im2col_3d, "im2col.3d") \
1182 M(im2col_4d, "im2col.4d") \
1183 M(im2col_5d, "im2col.5d") \
1184 M(im2col_w_3d, "im2col.w.3d") \
1185 M(im2col_w_4d, "im2col.w.4d") \
1186 M(im2col_w_5d, "im2col.w.5d") \
1187 M(im2col_w_128_3d, "im2col.w.128.3d") \
1188 M(im2col_w_128_4d, "im2col.w.128.4d") \
1189 M(im2col_w_128_5d, "im2col.w.128.5d")
1201 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1204#define G2S_ID(ID_SUFFIX, NAME) \
1205 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1215 size_t NumParams =
F->getFunctionType()->getNumParams();
1219 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1226 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1227 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1228 "expected the i64 cache-hint after the multicast mask");
1229 Type *MaskTy = Params[MaskIdx];
1244 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1247#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1248 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1260 if (!
F->getFunctionType()
1261 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1277 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1282 size_t NumParams =
F->getFunctionType()->getNumParams();
1283 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1287 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1292 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1305 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1310 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1313 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1333 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1336 auto [RedOpName, ShapeName] = Name.split(
'.');
1341 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1342 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1343 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1344 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1345 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1346 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1347 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1348 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1354 if (Name.consume_front(
"mapa.shared.cluster"))
1355 if (
F->getReturnType()->getPointerAddressSpace() ==
1357 return Intrinsic::nvvm_mapa_shared_cluster;
1359 if (Name.consume_front(
"cp.async.bulk.")) {
1362 .
Case(
"shared.cta.to.cluster",
1363 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1367 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1377 if (!Name.consume_front(
"tcgen05.commit."))
1380 if (Name.consume_front(
"shared."))
1382 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1383 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1386 if (Name.consume_front(
"mc.shared.")) {
1388 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1392 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1393 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1402 if (
F->arg_size() != 2)
1405 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1406 Name.consume_front(
"tcgen05.alloc."))
1408 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1409 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1412 if (Name.consume_front(
"tcgen05.dealloc."))
1414 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1415 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1422 if (Name.consume_front(
"fma.rn."))
1424 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1425 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1426 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1427 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1430 if (Name.consume_front(
"fmax."))
1432 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1433 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1434 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1435 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1436 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1437 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1438 .
Case(
"ftz.nan.xorsign.abs.bf16",
1439 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1440 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1441 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1442 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1443 .
Case(
"ftz.xorsign.abs.bf16x2",
1444 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1445 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1446 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1447 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1448 .
Case(
"nan.xorsign.abs.bf16x2",
1449 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1450 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1451 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1454 if (Name.consume_front(
"fmin."))
1456 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1457 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1458 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1459 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1460 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1461 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1462 .
Case(
"ftz.nan.xorsign.abs.bf16",
1463 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1464 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1465 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1466 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1467 .
Case(
"ftz.xorsign.abs.bf16x2",
1468 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1469 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1470 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1471 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1472 .
Case(
"nan.xorsign.abs.bf16x2",
1473 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1474 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1475 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1478 if (Name.consume_front(
"neg."))
1480 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1481 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1490 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1495 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1498 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1501 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1502 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1508static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1510 auto [Modifiers,
Type] = Name.rsplit(
'.');
1512 return std::nullopt;
1522 return std::nullopt;
1525 .
Case(
"", Intrinsic::nvvm_fadd)
1526 .
Case(
".ftz", Intrinsic::nvvm_fadd_ftz)
1527 .
Case(
".sat", Intrinsic::nvvm_fadd_sat)
1528 .
Case(
".ftz.sat", Intrinsic::nvvm_fadd_ftz_sat)
1531 return std::nullopt;
1537 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1540 return Intrinsic::nvvm_mbarrier_init;
1544 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1545 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1546 Name.consume_front(
"param");
1550 if (!Name.consume_front(
"vp."))
1579 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1580 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1587 if (!Name.consume_front(
"vp."))
1607 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1608 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1613 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1625 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1636 if (Name.starts_with(
"to.fp16")) {
1640 FuncTy->getReturnType());
1643 if (Name.starts_with(
"from.fp16")) {
1647 FuncTy->getReturnType());
1656 if (Defaults.empty())
1662 unsigned FullArgCount = FirstDefault + Defaults.size();
1665 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1668 return FullArgCount;
1675 if (FullArgCount == 0)
1681 "total number of default args does not match intrinsic signature");
1686 bool CanUpgradeDebugIntrinsicsToRecords) {
1687 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1692 if (!Name.consume_front(
"llvm.") || Name.empty())
1698 bool IsArm = Name.consume_front(
"arm.");
1699 if (IsArm || Name.consume_front(
"aarch64.")) {
1705 if (Name.consume_front(
"amdgcn.")) {
1706 if (Name ==
"alignbit") {
1709 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1713 if (Name.consume_front(
"atomic.")) {
1714 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1715 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1724 if (Name.starts_with(
"addrspacecast.nonnull")) {
1731 switch (
F->getIntrinsicID()) {
1735 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1736 if (
F->arg_size() == 7) {
1741 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1742 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1743 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1744 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1745 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1746 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1747 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1748 if (
F->arg_size() == 8) {
1755 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1756 Name.consume_front(
"flat.atomic.")) {
1757 if (Name.starts_with(
"fadd") ||
1759 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1760 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1768 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1773 if (Name.starts_with(
"ldexp.")) {
1776 F->getParent(), Intrinsic::ldexp,
1777 {F->getReturnType(), F->getArg(1)->getType()});
1786 if (
F->arg_size() == 1) {
1787 if (Name.consume_front(
"convert.")) {
1801 F->arg_begin()->getType());
1807 if (Name ==
"coro.end" &&
1808 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1809 CoroEndID = Intrinsic::coro_end;
1810 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1811 CoroEndID = Intrinsic::coro_end_async;
1822 if (Name.consume_front(
"dbg.")) {
1824 if (CanUpgradeDebugIntrinsicsToRecords) {
1825 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1826 Name ==
"declare" || Name ==
"label") {
1835 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1838 Intrinsic::dbg_value);
1845 if (Name.consume_front(
"experimental.vector.")) {
1851 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1852 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1853 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1854 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1855 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1857 Intrinsic::vector_partial_reduce_add)
1860 const auto *FT =
F->getFunctionType();
1862 if (ID == Intrinsic::vector_extract ||
1863 ID == Intrinsic::vector_interleave2)
1866 if (ID != Intrinsic::vector_interleave2)
1868 if (ID == Intrinsic::vector_insert ||
1869 ID == Intrinsic::vector_partial_reduce_add)
1877 if (Name.consume_front(
"reduce.")) {
1879 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1880 if (R.match(Name, &
Groups))
1882 .
Case(
"add", Intrinsic::vector_reduce_add)
1883 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1884 .
Case(
"and", Intrinsic::vector_reduce_and)
1885 .
Case(
"or", Intrinsic::vector_reduce_or)
1886 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1887 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1888 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1889 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1890 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1891 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1892 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1897 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1902 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1903 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1908 auto Args =
F->getFunctionType()->params();
1910 {Args[V2 ? 1 : 0]});
1916 if (Name.consume_front(
"splice"))
1920 if (Name.consume_front(
"experimental.stepvector.")) {
1924 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1929 if (Name.starts_with(
"flt.rounds")) {
1932 Intrinsic::get_rounding);
1937 if (Name.starts_with(
"invariant.group.barrier")) {
1939 auto Args =
F->getFunctionType()->params();
1940 Type* ObjectPtr[1] = {Args[0]};
1943 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1948 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1949 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1950 if (IsLifetimeStart || IsLifetimeEnd) {
1951 if (
F->arg_size() == 2) {
1952 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1953 : Intrinsic::lifetime_end;
1958 F->getArg(1)->getType());
1960 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1980 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1981 .StartsWith(
"memmove.", Intrinsic::memmove)
1983 if (
F->arg_size() == 5) {
1987 F->getFunctionType()->params().slice(0, 3);
1993 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1996 const auto *FT =
F->getFunctionType();
1997 Type *ParamTypes[2] = {
1998 FT->getParamType(0),
2002 Intrinsic::memset, ParamTypes);
2008 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2009 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2010 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2011 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2013 if (MaskedID &&
F->arg_size() == 4) {
2015 if (MaskedID == Intrinsic::masked_load ||
2016 MaskedID == Intrinsic::masked_gather) {
2018 F->getParent(), MaskedID,
2019 {F->getReturnType(), F->getArg(0)->getType()});
2023 F->getParent(), MaskedID,
2024 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2030 if (Name.consume_front(
"nvvm.")) {
2032 if (
F->arg_size() == 1) {
2035 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2036 .Case(
"clz.i", Intrinsic::ctlz)
2037 .
Case(
"popc.i", Intrinsic::ctpop)
2041 {F->getReturnType()});
2044 }
else if (
F->arg_size() == 2) {
2047 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2048 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2049 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2050 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2054 {F->getReturnType()});
2091 F->getParent(), IID,
F->getReturnType(),
2092 F->getFunctionType()->params());
2103 {F->getArg(0)->getType()});
2151 F->getArg(0)->getType());
2159 bool Expand =
false;
2160 if (Name.consume_front(
"abs."))
2163 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2164 else if (Name.consume_front(
"fabs."))
2166 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2167 else if (Name.consume_front(
"add."))
2170 else if (Name.consume_front(
"ex2.approx."))
2173 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2174 else if (Name.consume_front(
"atomic.load."))
2183 else if (Name.consume_front(
"atomic."))
2198 else if (Name.consume_front(
"bitcast."))
2201 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2202 else if (Name.consume_front(
"rotate."))
2204 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2205 else if (Name.consume_front(
"ptr.gen.to."))
2208 else if (Name.consume_front(
"ptr."))
2211 else if (Name.consume_front(
"ldg.global."))
2213 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2214 Name.starts_with(
"p."));
2217 .
Case(
"barrier0",
true)
2218 .
Case(
"barrier.n",
true)
2219 .
Case(
"barrier.sync.cnt",
true)
2220 .
Case(
"barrier.sync",
true)
2221 .
Case(
"barrier",
true)
2222 .
Case(
"bar.sync",
true)
2223 .
Case(
"barrier0.popc",
true)
2224 .
Case(
"barrier0.and",
true)
2225 .
Case(
"barrier0.or",
true)
2226 .
Case(
"clz.ll",
true)
2227 .
Case(
"popc.ll",
true)
2229 .
Case(
"swap.lo.hi.b64",
true)
2230 .
Case(
"tanh.approx.f32",
true)
2242 if (Name.starts_with(
"objectsize.")) {
2243 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2244 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2247 Intrinsic::objectsize, Tys);
2254 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2257 F->getParent(), Intrinsic::ptr_annotation,
2258 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2264 if (Name.consume_front(
"riscv.")) {
2267 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2268 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2269 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2270 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2273 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2286 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2287 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2296 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2297 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2298 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2299 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2304 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2313 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2315 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2324 if (Name ==
"stackprotectorcheck") {
2331 if (Name ==
"thread.pointer") {
2333 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2339 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2342 F->getParent(), Intrinsic::var_annotation,
2343 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2346 if (Name.consume_front(
"vector.splice")) {
2347 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2357 if (Name.consume_front(
"wasm.")) {
2360 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2361 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2362 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2367 F->getReturnType());
2371 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2373 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2375 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2394 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2404 std::string
Name =
F->getName().str();
2407 Name,
F->getParent());
2418 if (Result != std::nullopt) {
2435 bool CanUpgradeDebugIntrinsicsToRecords) {
2455 GV->
getName() ==
"llvm.global_dtors")) ||
2470 unsigned N =
Init->getNumOperands();
2471 std::vector<Constant *> NewCtors(
N);
2472 for (
unsigned i = 0; i !=
N; ++i) {
2475 Ctor->getAggregateElement(1),
2489 unsigned NumElts = ResultTy->getNumElements() * 8;
2493 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2503 for (
unsigned l = 0; l != NumElts; l += 16)
2504 for (
unsigned i = 0; i != 16; ++i) {
2505 unsigned Idx = NumElts + i - Shift;
2507 Idx -= NumElts - 16;
2508 Idxs[l + i] = Idx + l;
2511 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2515 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2523 unsigned NumElts = ResultTy->getNumElements() * 8;
2527 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2537 for (
unsigned l = 0; l != NumElts; l += 16)
2538 for (
unsigned i = 0; i != 16; ++i) {
2539 unsigned Idx = i + Shift;
2541 Idx += NumElts - 16;
2542 Idxs[l + i] = Idx + l;
2545 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2549 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2557 Mask = Builder.CreateBitCast(Mask, MaskTy);
2563 for (
unsigned i = 0; i != NumElts; ++i)
2565 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2576 if (
C->isAllOnesValue())
2581 return Builder.CreateSelect(Mask, Op0, Op1);
2588 if (
C->isAllOnesValue())
2592 Mask->getType()->getIntegerBitWidth());
2593 Mask = Builder.CreateBitCast(Mask, MaskTy);
2594 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2595 return Builder.CreateSelect(Mask, Op0, Op1);
2608 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2609 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2614 ShiftVal &= (NumElts - 1);
2623 if (ShiftVal > 16) {
2631 for (
unsigned l = 0; l < NumElts; l += 16) {
2632 for (
unsigned i = 0; i != 16; ++i) {
2633 unsigned Idx = ShiftVal + i;
2634 if (!IsVALIGN && Idx >= 16)
2635 Idx += NumElts - 16;
2636 Indices[l + i] = Idx + l;
2641 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2647 bool ZeroMask,
bool IndexForm) {
2650 unsigned EltWidth = Ty->getScalarSizeInBits();
2651 bool IsFloat = Ty->isFPOrFPVectorTy();
2653 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2654 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2655 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2656 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2657 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2658 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2659 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2660 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2661 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2662 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2663 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2664 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2665 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2666 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2667 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2668 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2669 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2670 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2671 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2672 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2673 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2674 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2675 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2676 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2677 else if (VecWidth == 128 && EltWidth == 16)
2678 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2679 else if (VecWidth == 256 && EltWidth == 16)
2680 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2681 else if (VecWidth == 512 && EltWidth == 16)
2682 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2683 else if (VecWidth == 128 && EltWidth == 8)
2684 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2685 else if (VecWidth == 256 && EltWidth == 8)
2686 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2687 else if (VecWidth == 512 && EltWidth == 8)
2688 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2699 Value *V = Builder.CreateIntrinsic(IID, Args);
2711 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2722 bool IsRotateRight) {
2732 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2733 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2736 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2737 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2782 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2787 bool IsShiftRight,
bool ZeroMask) {
2801 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2802 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2805 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2806 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2821 const Align Alignment =
2823 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2828 if (
C->isAllOnesValue())
2829 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2834 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2840 const Align Alignment =
2849 if (
C->isAllOnesValue())
2850 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2855 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2861 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2862 {Op0, Builder.getInt1(
false)});
2877 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2878 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2879 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2880 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2881 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2884 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2885 LHS = Builder.CreateAnd(
LHS, Mask);
2886 RHS = Builder.CreateAnd(
RHS, Mask);
2903 if (!
C || !
C->isAllOnesValue())
2904 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2909 for (
unsigned i = 0; i != NumElts; ++i)
2911 for (
unsigned i = NumElts; i != 8; ++i)
2912 Indices[i] = NumElts + i % NumElts;
2913 Vec = Builder.CreateShuffleVector(Vec,
2917 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2921 unsigned CC,
bool Signed) {
2929 }
else if (CC == 7) {
2965 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2966 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2968 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2969 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2978 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2984 Name = Name.substr(12);
2989 if (Name.starts_with(
"max.p")) {
2990 if (VecWidth == 128 && EltWidth == 32)
2991 IID = Intrinsic::x86_sse_max_ps;
2992 else if (VecWidth == 128 && EltWidth == 64)
2993 IID = Intrinsic::x86_sse2_max_pd;
2994 else if (VecWidth == 256 && EltWidth == 32)
2995 IID = Intrinsic::x86_avx_max_ps_256;
2996 else if (VecWidth == 256 && EltWidth == 64)
2997 IID = Intrinsic::x86_avx_max_pd_256;
3000 }
else if (Name.starts_with(
"min.p")) {
3001 if (VecWidth == 128 && EltWidth == 32)
3002 IID = Intrinsic::x86_sse_min_ps;
3003 else if (VecWidth == 128 && EltWidth == 64)
3004 IID = Intrinsic::x86_sse2_min_pd;
3005 else if (VecWidth == 256 && EltWidth == 32)
3006 IID = Intrinsic::x86_avx_min_ps_256;
3007 else if (VecWidth == 256 && EltWidth == 64)
3008 IID = Intrinsic::x86_avx_min_pd_256;
3011 }
else if (Name.starts_with(
"pshuf.b.")) {
3012 if (VecWidth == 128)
3013 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3014 else if (VecWidth == 256)
3015 IID = Intrinsic::x86_avx2_pshuf_b;
3016 else if (VecWidth == 512)
3017 IID = Intrinsic::x86_avx512_pshuf_b_512;
3020 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3021 if (VecWidth == 128)
3022 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3023 else if (VecWidth == 256)
3024 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3025 else if (VecWidth == 512)
3026 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3029 }
else if (Name.starts_with(
"pmulh.w.")) {
3030 if (VecWidth == 128)
3031 IID = Intrinsic::x86_sse2_pmulh_w;
3032 else if (VecWidth == 256)
3033 IID = Intrinsic::x86_avx2_pmulh_w;
3034 else if (VecWidth == 512)
3035 IID = Intrinsic::x86_avx512_pmulh_w_512;
3038 }
else if (Name.starts_with(
"pmulhu.w.")) {
3039 if (VecWidth == 128)
3040 IID = Intrinsic::x86_sse2_pmulhu_w;
3041 else if (VecWidth == 256)
3042 IID = Intrinsic::x86_avx2_pmulhu_w;
3043 else if (VecWidth == 512)
3044 IID = Intrinsic::x86_avx512_pmulhu_w_512;
3047 }
else if (Name.starts_with(
"pmaddw.d.")) {
3048 if (VecWidth == 128)
3049 IID = Intrinsic::x86_sse2_pmadd_wd;
3050 else if (VecWidth == 256)
3051 IID = Intrinsic::x86_avx2_pmadd_wd;
3052 else if (VecWidth == 512)
3053 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3056 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3057 if (VecWidth == 128)
3058 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3059 else if (VecWidth == 256)
3060 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3061 else if (VecWidth == 512)
3062 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3065 }
else if (Name.starts_with(
"packsswb.")) {
3066 if (VecWidth == 128)
3067 IID = Intrinsic::x86_sse2_packsswb_128;
3068 else if (VecWidth == 256)
3069 IID = Intrinsic::x86_avx2_packsswb;
3070 else if (VecWidth == 512)
3071 IID = Intrinsic::x86_avx512_packsswb_512;
3074 }
else if (Name.starts_with(
"packssdw.")) {
3075 if (VecWidth == 128)
3076 IID = Intrinsic::x86_sse2_packssdw_128;
3077 else if (VecWidth == 256)
3078 IID = Intrinsic::x86_avx2_packssdw;
3079 else if (VecWidth == 512)
3080 IID = Intrinsic::x86_avx512_packssdw_512;
3083 }
else if (Name.starts_with(
"packuswb.")) {
3084 if (VecWidth == 128)
3085 IID = Intrinsic::x86_sse2_packuswb_128;
3086 else if (VecWidth == 256)
3087 IID = Intrinsic::x86_avx2_packuswb;
3088 else if (VecWidth == 512)
3089 IID = Intrinsic::x86_avx512_packuswb_512;
3092 }
else if (Name.starts_with(
"packusdw.")) {
3093 if (VecWidth == 128)
3094 IID = Intrinsic::x86_sse41_packusdw;
3095 else if (VecWidth == 256)
3096 IID = Intrinsic::x86_avx2_packusdw;
3097 else if (VecWidth == 512)
3098 IID = Intrinsic::x86_avx512_packusdw_512;
3101 }
else if (Name.starts_with(
"vpermilvar.")) {
3102 if (VecWidth == 128 && EltWidth == 32)
3103 IID = Intrinsic::x86_avx_vpermilvar_ps;
3104 else if (VecWidth == 128 && EltWidth == 64)
3105 IID = Intrinsic::x86_avx_vpermilvar_pd;
3106 else if (VecWidth == 256 && EltWidth == 32)
3107 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3108 else if (VecWidth == 256 && EltWidth == 64)
3109 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3110 else if (VecWidth == 512 && EltWidth == 32)
3111 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3112 else if (VecWidth == 512 && EltWidth == 64)
3113 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3116 }
else if (Name ==
"cvtpd2dq.256") {
3117 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3118 }
else if (Name ==
"cvtpd2ps.256") {
3119 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3120 }
else if (Name ==
"cvttpd2dq.256") {
3121 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3122 }
else if (Name ==
"cvttps2dq.128") {
3123 IID = Intrinsic::x86_sse2_cvttps2dq;
3124 }
else if (Name ==
"cvttps2dq.256") {
3125 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3126 }
else if (Name.starts_with(
"permvar.")) {
3128 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3129 IID = Intrinsic::x86_avx2_permps;
3130 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3131 IID = Intrinsic::x86_avx2_permd;
3132 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3133 IID = Intrinsic::x86_avx512_permvar_df_256;
3134 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3135 IID = Intrinsic::x86_avx512_permvar_di_256;
3136 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3137 IID = Intrinsic::x86_avx512_permvar_sf_512;
3138 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3139 IID = Intrinsic::x86_avx512_permvar_si_512;
3140 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3141 IID = Intrinsic::x86_avx512_permvar_df_512;
3142 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3143 IID = Intrinsic::x86_avx512_permvar_di_512;
3144 else if (VecWidth == 128 && EltWidth == 16)
3145 IID = Intrinsic::x86_avx512_permvar_hi_128;
3146 else if (VecWidth == 256 && EltWidth == 16)
3147 IID = Intrinsic::x86_avx512_permvar_hi_256;
3148 else if (VecWidth == 512 && EltWidth == 16)
3149 IID = Intrinsic::x86_avx512_permvar_hi_512;
3150 else if (VecWidth == 128 && EltWidth == 8)
3151 IID = Intrinsic::x86_avx512_permvar_qi_128;
3152 else if (VecWidth == 256 && EltWidth == 8)
3153 IID = Intrinsic::x86_avx512_permvar_qi_256;
3154 else if (VecWidth == 512 && EltWidth == 8)
3155 IID = Intrinsic::x86_avx512_permvar_qi_512;
3158 }
else if (Name.starts_with(
"dbpsadbw.")) {
3159 if (VecWidth == 128)
3160 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3161 else if (VecWidth == 256)
3162 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3163 else if (VecWidth == 512)
3164 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3167 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3168 if (VecWidth == 128)
3169 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3170 else if (VecWidth == 256)
3171 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3172 else if (VecWidth == 512)
3173 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3176 }
else if (Name.starts_with(
"conflict.")) {
3177 if (Name[9] ==
'd' && VecWidth == 128)
3178 IID = Intrinsic::x86_avx512_conflict_d_128;
3179 else if (Name[9] ==
'd' && VecWidth == 256)
3180 IID = Intrinsic::x86_avx512_conflict_d_256;
3181 else if (Name[9] ==
'd' && VecWidth == 512)
3182 IID = Intrinsic::x86_avx512_conflict_d_512;
3183 else if (Name[9] ==
'q' && VecWidth == 128)
3184 IID = Intrinsic::x86_avx512_conflict_q_128;
3185 else if (Name[9] ==
'q' && VecWidth == 256)
3186 IID = Intrinsic::x86_avx512_conflict_q_256;
3187 else if (Name[9] ==
'q' && VecWidth == 512)
3188 IID = Intrinsic::x86_avx512_conflict_q_512;
3191 }
else if (Name.starts_with(
"pavg.")) {
3192 if (Name[5] ==
'b' && VecWidth == 128)
3193 IID = Intrinsic::x86_sse2_pavg_b;
3194 else if (Name[5] ==
'b' && VecWidth == 256)
3195 IID = Intrinsic::x86_avx2_pavg_b;
3196 else if (Name[5] ==
'b' && VecWidth == 512)
3197 IID = Intrinsic::x86_avx512_pavg_b_512;
3198 else if (Name[5] ==
'w' && VecWidth == 128)
3199 IID = Intrinsic::x86_sse2_pavg_w;
3200 else if (Name[5] ==
'w' && VecWidth == 256)
3201 IID = Intrinsic::x86_avx2_pavg_w;
3202 else if (Name[5] ==
'w' && VecWidth == 512)
3203 IID = Intrinsic::x86_avx512_pavg_w_512;
3212 Rep = Builder.CreateIntrinsic(IID, Args);
3223 if (AsmStr->find(
"mov\tfp") == 0 &&
3224 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3225 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3226 AsmStr->replace(Pos, 1,
";");
3232 Value *Rep =
nullptr;
3234 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3236 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3237 {Arg, Builder.getTrue()},
3239 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3240 Type *Ty = (Name ==
"abs.bf16")
3244 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3245 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3246 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3247 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3248 : Intrinsic::nvvm_fabs;
3249 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3250 }
else if (Name.consume_front(
"add.")) {
3253 assert(
FAdd &&
"unsupported nvvm.add.* intrinsic");
3256 Rep = Builder.CreateIntrinsic(
3258 {A, CI->getArgOperand(1),
3259 Builder.getInt32(static_cast<int>(RoundingMode))});
3260 }
else if (Name.consume_front(
"ex2.approx.")) {
3262 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3263 : Intrinsic::nvvm_ex2_approx;
3264 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3265 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3266 Name.starts_with(
"atomic.load.add.f64.p")) {
3269 Rep = Builder.CreateAtomicRMW(
3275 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3276 Name.starts_with(
"atomic.load.dec.32.p")) {
3281 Rep = Builder.CreateAtomicRMW(
3285 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3291 Op.contains(
".cta.") ?
"block" :
"");
3292 if (
Op.starts_with(
"cas.")) {
3294 Value *Pair = Builder.CreateAtomicCmpXchg(
3297 Rep = Builder.CreateExtractValue(Pair, 0);
3315 "unexpected nvvm scoped atomic intrinsic");
3316 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3319 }
else if (Name ==
"clz.ll") {
3322 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3323 {Arg, Builder.getFalse()},
3325 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3326 }
else if (Name ==
"popc.ll") {
3330 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3331 Arg,
nullptr,
"ctpop");
3332 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3333 }
else if (Name ==
"h2f") {
3335 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3336 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3337 }
else if (Name.consume_front(
"bitcast.") &&
3338 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3341 }
else if (Name ==
"rotate.b32") {
3344 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3345 {Arg, Arg, ShiftAmt});
3346 }
else if (Name ==
"rotate.b64") {
3350 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3351 {Arg, Arg, ZExtShiftAmt});
3352 }
else if (Name ==
"rotate.right.b64") {
3356 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3357 {Arg, Arg, ZExtShiftAmt});
3358 }
else if (Name ==
"swap.lo.hi.b64") {
3361 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3362 {Arg, Arg, Builder.getInt64(32)});
3363 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3366 Name.starts_with(
".to.gen"))) {
3368 }
else if (Name.consume_front(
"ldg.global")) {
3372 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3375 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3377 }
else if (Name ==
"tanh.approx.f32") {
3381 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3383 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3385 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3386 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3388 }
else if (Name ==
"barrier") {
3389 Rep = Builder.CreateIntrinsic(
3390 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3392 }
else if (Name ==
"barrier.sync") {
3393 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3395 }
else if (Name ==
"barrier.sync.cnt") {
3396 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3398 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3399 Name ==
"barrier0.or") {
3401 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3405 .
Case(
"barrier0.popc",
3406 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3407 .
Case(
"barrier0.and",
3408 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3409 .
Case(
"barrier0.or",
3410 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3411 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3412 Rep = Builder.CreateZExt(Bar, CI->
getType());
3426 ? Builder.CreateBitCast(Arg, NewType)
3429 Rep = Builder.CreateCall(NewFn, Args);
3430 if (
F->getReturnType()->isIntegerTy())
3431 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3441 Value *Rep =
nullptr;
3443 if (Name.starts_with(
"sse4a.movnt.")) {
3455 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3458 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3459 }
else if (Name.starts_with(
"avx.movnt.") ||
3460 Name.starts_with(
"avx512.storent.")) {
3472 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3473 }
else if (Name ==
"sse2.storel.dq") {
3478 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3479 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3480 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3481 }
else if (Name.starts_with(
"sse.storeu.") ||
3482 Name.starts_with(
"sse2.storeu.") ||
3483 Name.starts_with(
"avx.storeu.")) {
3486 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3487 }
else if (Name ==
"avx512.mask.store.ss") {
3491 }
else if (Name.starts_with(
"avx512.mask.store")) {
3493 bool Aligned = Name[17] !=
'u';
3496 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3499 bool CmpEq = Name[9] ==
'e';
3502 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3503 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3510 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3511 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3513 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3514 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3515 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3516 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3517 Name.starts_with(
"sse2.sqrt.p") ||
3518 Name.starts_with(
"sse.sqrt.p")) {
3519 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3520 {CI->getArgOperand(0)});
3521 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3525 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3526 : Intrinsic::x86_avx512_sqrt_pd_512;
3529 Rep = Builder.CreateIntrinsic(IID, Args);
3531 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3532 {CI->getArgOperand(0)});
3536 }
else if (Name.starts_with(
"avx512.ptestm") ||
3537 Name.starts_with(
"avx512.ptestnm")) {
3541 Rep = Builder.CreateAnd(Op0, Op1);
3547 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3549 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3552 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3555 }
else if (Name.starts_with(
"avx512.kunpck")) {
3560 for (
unsigned i = 0; i != NumElts; ++i)
3569 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3570 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3571 }
else if (Name ==
"avx512.kand.w") {
3574 Rep = Builder.CreateAnd(
LHS,
RHS);
3575 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3576 }
else if (Name ==
"avx512.kandn.w") {
3579 LHS = Builder.CreateNot(
LHS);
3580 Rep = Builder.CreateAnd(
LHS,
RHS);
3581 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3582 }
else if (Name ==
"avx512.kor.w") {
3585 Rep = Builder.CreateOr(
LHS,
RHS);
3586 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3587 }
else if (Name ==
"avx512.kxor.w") {
3590 Rep = Builder.CreateXor(
LHS,
RHS);
3591 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3592 }
else if (Name ==
"avx512.kxnor.w") {
3595 LHS = Builder.CreateNot(
LHS);
3596 Rep = Builder.CreateXor(
LHS,
RHS);
3597 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3598 }
else if (Name ==
"avx512.knot.w") {
3600 Rep = Builder.CreateNot(Rep);
3601 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3602 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3605 Rep = Builder.CreateOr(
LHS,
RHS);
3606 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3608 if (Name[14] ==
'c')
3612 Rep = Builder.CreateICmpEQ(Rep,
C);
3613 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3614 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3615 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3616 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3617 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3620 ConstantInt::get(I32Ty, 0));
3622 ConstantInt::get(I32Ty, 0));
3624 if (Name.contains(
".add."))
3625 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3626 else if (Name.contains(
".sub."))
3627 EltOp = Builder.CreateFSub(Elt0, Elt1);
3628 else if (Name.contains(
".mul."))
3629 EltOp = Builder.CreateFMul(Elt0, Elt1);
3631 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3632 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3633 ConstantInt::get(I32Ty, 0));
3634 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3636 bool CmpEq = Name[16] ==
'e';
3638 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3640 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3647 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3650 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3653 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3660 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3662 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3663 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3665 if (VecWidth == 128 && EltWidth == 32)
3666 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3667 else if (VecWidth == 256 && EltWidth == 32)
3668 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3669 else if (VecWidth == 512 && EltWidth == 32)
3670 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3671 else if (VecWidth == 128 && EltWidth == 64)
3672 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3673 else if (VecWidth == 256 && EltWidth == 64)
3674 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3675 else if (VecWidth == 512 && EltWidth == 64)
3676 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3683 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3686 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3687 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3689 if (VecWidth == 128 && EltWidth == 32)
3690 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3691 else if (VecWidth == 256 && EltWidth == 32)
3692 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3693 else if (VecWidth == 512 && EltWidth == 32)
3694 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3695 else if (VecWidth == 128 && EltWidth == 64)
3696 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3697 else if (VecWidth == 256 && EltWidth == 64)
3698 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3699 else if (VecWidth == 512 && EltWidth == 64)
3700 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3705 if (VecWidth == 512)
3707 Args.push_back(Mask);
3709 Rep = Builder.CreateIntrinsic(IID, Args);
3710 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3714 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3717 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3718 Name.starts_with(
"avx512.cvtw2mask.") ||
3719 Name.starts_with(
"avx512.cvtd2mask.") ||
3720 Name.starts_with(
"avx512.cvtq2mask.")) {
3725 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3726 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3727 Name.starts_with(
"avx512.mask.pabs")) {
3729 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3730 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3731 Name.starts_with(
"avx512.mask.pmaxs")) {
3733 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3734 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3735 Name.starts_with(
"avx512.mask.pmaxu")) {
3737 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3738 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3739 Name.starts_with(
"avx512.mask.pmins")) {
3741 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3742 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3743 Name.starts_with(
"avx512.mask.pminu")) {
3745 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3746 Name ==
"avx512.pmulu.dq.512" ||
3747 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3749 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3750 Name ==
"avx512.pmul.dq.512" ||
3751 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3753 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3754 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3759 }
else if (Name ==
"avx512.cvtusi2sd") {
3764 }
else if (Name ==
"sse2.cvtss2sd") {
3766 Rep = Builder.CreateFPExt(
3769 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3770 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3771 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3772 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3773 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3774 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3775 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3776 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3777 Name ==
"avx512.mask.cvtqq2ps.256" ||
3778 Name ==
"avx512.mask.cvtqq2ps.512" ||
3779 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3780 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3781 Name ==
"avx.cvt.ps2.pd.256" ||
3782 Name ==
"avx512.mask.cvtps2pd.128" ||
3783 Name ==
"avx512.mask.cvtps2pd.256") {
3788 unsigned NumDstElts = DstTy->getNumElements();
3789 if (NumDstElts < SrcTy->getNumElements()) {
3790 assert(NumDstElts == 2 &&
"Unexpected vector size");
3791 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3794 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3795 bool IsUnsigned = Name.contains(
"cvtu");
3797 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3801 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3802 : Intrinsic::x86_avx512_sitofp_round;
3803 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3806 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3807 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3813 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3814 Name.starts_with(
"vcvtph2ps.")) {
3818 unsigned NumDstElts = DstTy->getNumElements();
3819 if (NumDstElts != SrcTy->getNumElements()) {
3820 assert(NumDstElts == 4 &&
"Unexpected vector size");
3821 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3823 Rep = Builder.CreateBitCast(
3825 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3829 }
else if (Name.starts_with(
"avx512.mask.load")) {
3831 bool Aligned = Name[16] !=
'u';
3834 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3838 ResultTy->getNumElements());
3839 Rep = Builder.CreateIntrinsic(
3840 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3842 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3848 Rep = Builder.CreateIntrinsic(
3849 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3851 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3852 Name.starts_with(
"avx512.mask.expand.")) {
3856 ResultTy->getNumElements());
3858 bool IsCompress = Name[12] ==
'c';
3859 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3860 : Intrinsic::x86_avx512_mask_expand;
3861 Rep = Builder.CreateIntrinsic(
3863 }
else if (Name.starts_with(
"xop.vpcom")) {
3865 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3866 Name.ends_with(
"uq"))
3868 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3869 Name.ends_with(
"d") || Name.ends_with(
"q"))
3878 Name = Name.substr(9);
3879 if (Name.starts_with(
"lt"))
3881 else if (Name.starts_with(
"le"))
3883 else if (Name.starts_with(
"gt"))
3885 else if (Name.starts_with(
"ge"))
3887 else if (Name.starts_with(
"eq"))
3889 else if (Name.starts_with(
"ne"))
3891 else if (Name.starts_with(
"false"))
3893 else if (Name.starts_with(
"true"))
3900 }
else if (Name.starts_with(
"xop.vpcmov")) {
3902 Value *NotSel = Builder.CreateNot(Sel);
3905 Rep = Builder.CreateOr(Sel0, Sel1);
3906 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3907 Name.starts_with(
"avx512.mask.prol")) {
3909 }
else if (Name.starts_with(
"avx512.pror") ||
3910 Name.starts_with(
"avx512.mask.pror")) {
3912 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3913 Name.starts_with(
"avx512.mask.vpshld") ||
3914 Name.starts_with(
"avx512.maskz.vpshld")) {
3915 bool ZeroMask = Name[11] ==
'z';
3917 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3918 Name.starts_with(
"avx512.mask.vpshrd") ||
3919 Name.starts_with(
"avx512.maskz.vpshrd")) {
3920 bool ZeroMask = Name[11] ==
'z';
3922 }
else if (Name ==
"sse42.crc32.64.8") {
3925 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3927 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3928 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3929 Name.starts_with(
"avx512.vbroadcast.s")) {
3932 Type *EltTy = VecTy->getElementType();
3933 unsigned EltNum = VecTy->getNumElements();
3937 for (
unsigned I = 0;
I < EltNum; ++
I)
3938 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3939 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3940 Name.starts_with(
"sse41.pmovzx") ||
3941 Name.starts_with(
"avx2.pmovsx") ||
3942 Name.starts_with(
"avx2.pmovzx") ||
3943 Name.starts_with(
"avx512.mask.pmovsx") ||
3944 Name.starts_with(
"avx512.mask.pmovzx")) {
3946 unsigned NumDstElts = DstTy->getNumElements();
3950 for (
unsigned i = 0; i != NumDstElts; ++i)
3955 bool DoSext = Name.contains(
"pmovsx");
3957 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3962 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3963 Name ==
"avx512.mask.pmov.qd.512" ||
3964 Name ==
"avx512.mask.pmov.wb.256" ||
3965 Name ==
"avx512.mask.pmov.wb.512") {
3970 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3971 Name ==
"avx2.vbroadcasti128") {
3977 if (NumSrcElts == 2)
3980 Rep = Builder.CreateShuffleVector(
Load,
3982 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3983 Name.starts_with(
"avx512.mask.shuf.f")) {
3988 unsigned ControlBitsMask = NumLanes - 1;
3989 unsigned NumControlBits = NumLanes / 2;
3992 for (
unsigned l = 0; l != NumLanes; ++l) {
3993 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3995 if (l >= NumLanes / 2)
3996 LaneMask += NumLanes;
3997 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3998 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4004 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4005 Name.starts_with(
"avx512.mask.broadcasti")) {
4008 unsigned NumDstElts =
4012 for (
unsigned i = 0; i != NumDstElts; ++i)
4013 ShuffleMask[i] = i % NumSrcElts;
4019 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4020 Name.starts_with(
"avx2.vbroadcast") ||
4021 Name.starts_with(
"avx512.pbroadcast") ||
4022 Name.starts_with(
"avx512.mask.broadcast.s")) {
4029 Rep = Builder.CreateShuffleVector(
Op, M);
4034 }
else if (Name.starts_with(
"sse2.padds.") ||
4035 Name.starts_with(
"avx2.padds.") ||
4036 Name.starts_with(
"avx512.padds.") ||
4037 Name.starts_with(
"avx512.mask.padds.")) {
4039 }
else if (Name.starts_with(
"sse2.psubs.") ||
4040 Name.starts_with(
"avx2.psubs.") ||
4041 Name.starts_with(
"avx512.psubs.") ||
4042 Name.starts_with(
"avx512.mask.psubs.")) {
4044 }
else if (Name.starts_with(
"sse2.paddus.") ||
4045 Name.starts_with(
"avx2.paddus.") ||
4046 Name.starts_with(
"avx512.mask.paddus.")) {
4048 }
else if (Name.starts_with(
"sse2.psubus.") ||
4049 Name.starts_with(
"avx2.psubus.") ||
4050 Name.starts_with(
"avx512.mask.psubus.")) {
4052 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4057 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4061 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4066 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4071 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4072 Name ==
"avx512.psll.dq.512") {
4076 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4077 Name ==
"avx512.psrl.dq.512") {
4081 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4082 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4083 Name.starts_with(
"avx2.pblendd.")) {
4088 unsigned NumElts = VecTy->getNumElements();
4091 for (
unsigned i = 0; i != NumElts; ++i)
4092 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4094 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4095 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4096 Name ==
"avx2.vinserti128" ||
4097 Name.starts_with(
"avx512.mask.insert")) {
4101 unsigned DstNumElts =
4103 unsigned SrcNumElts =
4105 unsigned Scale = DstNumElts / SrcNumElts;
4112 for (
unsigned i = 0; i != SrcNumElts; ++i)
4114 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4115 Idxs[i] = SrcNumElts;
4116 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4130 for (
unsigned i = 0; i != DstNumElts; ++i)
4133 for (
unsigned i = 0; i != SrcNumElts; ++i)
4134 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4135 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4141 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4142 Name ==
"avx2.vextracti128" ||
4143 Name.starts_with(
"avx512.mask.vextract")) {
4146 unsigned DstNumElts =
4148 unsigned SrcNumElts =
4150 unsigned Scale = SrcNumElts / DstNumElts;
4157 for (
unsigned i = 0; i != DstNumElts; ++i) {
4158 Idxs[i] = i + (
Imm * DstNumElts);
4160 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4166 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4167 Name.starts_with(
"avx512.mask.perm.di.")) {
4171 unsigned NumElts = VecTy->getNumElements();
4174 for (
unsigned i = 0; i != NumElts; ++i)
4175 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4177 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4182 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4194 unsigned HalfSize = NumElts / 2;
4206 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4207 for (
unsigned i = 0; i < HalfSize; ++i)
4208 ShuffleMask[i] = StartIndex + i;
4211 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4212 for (
unsigned i = 0; i < HalfSize; ++i)
4213 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4215 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4217 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4218 Name.starts_with(
"avx512.mask.vpermil.p") ||
4219 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4223 unsigned NumElts = VecTy->getNumElements();
4225 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4226 unsigned IdxMask = ((1 << IdxSize) - 1);
4232 for (
unsigned i = 0; i != NumElts; ++i)
4233 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4235 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4240 }
else if (Name ==
"sse2.pshufl.w" ||
4241 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4246 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4250 for (
unsigned l = 0; l != NumElts; l += 8) {
4251 for (
unsigned i = 0; i != 4; ++i)
4252 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4253 for (
unsigned i = 4; i != 8; ++i)
4254 Idxs[i + l] = i + l;
4257 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4262 }
else if (Name ==
"sse2.pshufh.w" ||
4263 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4268 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4272 for (
unsigned l = 0; l != NumElts; l += 8) {
4273 for (
unsigned i = 0; i != 4; ++i)
4274 Idxs[i + l] = i + l;
4275 for (
unsigned i = 0; i != 4; ++i)
4276 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4279 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4284 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4291 unsigned HalfLaneElts = NumLaneElts / 2;
4294 for (
unsigned i = 0; i != NumElts; ++i) {
4296 Idxs[i] = i - (i % NumLaneElts);
4298 if ((i % NumLaneElts) >= HalfLaneElts)
4302 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4305 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4309 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4310 Name.starts_with(
"avx512.mask.movshdup") ||
4311 Name.starts_with(
"avx512.mask.movsldup")) {
4317 if (Name.starts_with(
"avx512.mask.movshdup."))
4321 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4322 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4323 Idxs[i + l + 0] = i + l +
Offset;
4324 Idxs[i + l + 1] = i + l +
Offset;
4327 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4331 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4332 Name.starts_with(
"avx512.mask.unpckl.")) {
4339 for (
int l = 0; l != NumElts; l += NumLaneElts)
4340 for (
int i = 0; i != NumLaneElts; ++i)
4341 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4343 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4347 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4348 Name.starts_with(
"avx512.mask.unpckh.")) {
4355 for (
int l = 0; l != NumElts; l += NumLaneElts)
4356 for (
int i = 0; i != NumLaneElts; ++i)
4357 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4359 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4363 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4364 Name.starts_with(
"avx512.mask.pand.")) {
4367 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4369 Rep = Builder.CreateBitCast(Rep, FTy);
4372 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4373 Name.starts_with(
"avx512.mask.pandn.")) {
4376 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4377 Rep = Builder.CreateAnd(Rep,
4379 Rep = Builder.CreateBitCast(Rep, FTy);
4382 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4383 Name.starts_with(
"avx512.mask.por.")) {
4386 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4388 Rep = Builder.CreateBitCast(Rep, FTy);
4391 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4392 Name.starts_with(
"avx512.mask.pxor.")) {
4395 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4397 Rep = Builder.CreateBitCast(Rep, FTy);
4400 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4404 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4408 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4412 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4413 if (Name.ends_with(
".512")) {
4415 if (Name[17] ==
's')
4416 IID = Intrinsic::x86_avx512_add_ps_512;
4418 IID = Intrinsic::x86_avx512_add_pd_512;
4420 Rep = Builder.CreateIntrinsic(
4428 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4429 if (Name.ends_with(
".512")) {
4431 if (Name[17] ==
's')
4432 IID = Intrinsic::x86_avx512_div_ps_512;
4434 IID = Intrinsic::x86_avx512_div_pd_512;
4436 Rep = Builder.CreateIntrinsic(
4444 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4445 if (Name.ends_with(
".512")) {
4447 if (Name[17] ==
's')
4448 IID = Intrinsic::x86_avx512_mul_ps_512;
4450 IID = Intrinsic::x86_avx512_mul_pd_512;
4452 Rep = Builder.CreateIntrinsic(
4460 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4461 if (Name.ends_with(
".512")) {
4463 if (Name[17] ==
's')
4464 IID = Intrinsic::x86_avx512_sub_ps_512;
4466 IID = Intrinsic::x86_avx512_sub_pd_512;
4468 Rep = Builder.CreateIntrinsic(
4476 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4477 Name.starts_with(
"avx512.mask.min.p")) &&
4478 Name.drop_front(18) ==
".512") {
4479 bool IsDouble = Name[17] ==
'd';
4480 bool IsMin = Name[13] ==
'i';
4482 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4483 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4486 Rep = Builder.CreateIntrinsic(
4491 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4493 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4494 {CI->getArgOperand(0), Builder.getInt1(false)});
4497 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4498 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4499 bool IsVariable = Name[16] ==
'v';
4500 char Size = Name[16] ==
'.' ? Name[17]
4501 : Name[17] ==
'.' ? Name[18]
4502 : Name[18] ==
'.' ? Name[19]
4506 if (IsVariable && Name[17] !=
'.') {
4507 if (
Size ==
'd' && Name[17] ==
'2')
4508 IID = Intrinsic::x86_avx2_psllv_q;
4509 else if (
Size ==
'd' && Name[17] ==
'4')
4510 IID = Intrinsic::x86_avx2_psllv_q_256;
4511 else if (
Size ==
's' && Name[17] ==
'4')
4512 IID = Intrinsic::x86_avx2_psllv_d;
4513 else if (
Size ==
's' && Name[17] ==
'8')
4514 IID = Intrinsic::x86_avx2_psllv_d_256;
4515 else if (
Size ==
'h' && Name[17] ==
'8')
4516 IID = Intrinsic::x86_avx512_psllv_w_128;
4517 else if (
Size ==
'h' && Name[17] ==
'1')
4518 IID = Intrinsic::x86_avx512_psllv_w_256;
4519 else if (Name[17] ==
'3' && Name[18] ==
'2')
4520 IID = Intrinsic::x86_avx512_psllv_w_512;
4523 }
else if (Name.ends_with(
".128")) {
4525 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4526 : Intrinsic::x86_sse2_psll_d;
4527 else if (
Size ==
'q')
4528 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4529 : Intrinsic::x86_sse2_psll_q;
4530 else if (
Size ==
'w')
4531 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4532 : Intrinsic::x86_sse2_psll_w;
4535 }
else if (Name.ends_with(
".256")) {
4537 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4538 : Intrinsic::x86_avx2_psll_d;
4539 else if (
Size ==
'q')
4540 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4541 : Intrinsic::x86_avx2_psll_q;
4542 else if (
Size ==
'w')
4543 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4544 : Intrinsic::x86_avx2_psll_w;
4549 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4550 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4551 : Intrinsic::x86_avx512_psll_d_512;
4552 else if (
Size ==
'q')
4553 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4554 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4555 : Intrinsic::x86_avx512_psll_q_512;
4556 else if (
Size ==
'w')
4557 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4558 : Intrinsic::x86_avx512_psll_w_512;
4564 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4565 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4566 bool IsVariable = Name[16] ==
'v';
4567 char Size = Name[16] ==
'.' ? Name[17]
4568 : Name[17] ==
'.' ? Name[18]
4569 : Name[18] ==
'.' ? Name[19]
4573 if (IsVariable && Name[17] !=
'.') {
4574 if (
Size ==
'd' && Name[17] ==
'2')
4575 IID = Intrinsic::x86_avx2_psrlv_q;
4576 else if (
Size ==
'd' && Name[17] ==
'4')
4577 IID = Intrinsic::x86_avx2_psrlv_q_256;
4578 else if (
Size ==
's' && Name[17] ==
'4')
4579 IID = Intrinsic::x86_avx2_psrlv_d;
4580 else if (
Size ==
's' && Name[17] ==
'8')
4581 IID = Intrinsic::x86_avx2_psrlv_d_256;
4582 else if (
Size ==
'h' && Name[17] ==
'8')
4583 IID = Intrinsic::x86_avx512_psrlv_w_128;
4584 else if (
Size ==
'h' && Name[17] ==
'1')
4585 IID = Intrinsic::x86_avx512_psrlv_w_256;
4586 else if (Name[17] ==
'3' && Name[18] ==
'2')
4587 IID = Intrinsic::x86_avx512_psrlv_w_512;
4590 }
else if (Name.ends_with(
".128")) {
4592 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4593 : Intrinsic::x86_sse2_psrl_d;
4594 else if (
Size ==
'q')
4595 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4596 : Intrinsic::x86_sse2_psrl_q;
4597 else if (
Size ==
'w')
4598 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4599 : Intrinsic::x86_sse2_psrl_w;
4602 }
else if (Name.ends_with(
".256")) {
4604 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4605 : Intrinsic::x86_avx2_psrl_d;
4606 else if (
Size ==
'q')
4607 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4608 : Intrinsic::x86_avx2_psrl_q;
4609 else if (
Size ==
'w')
4610 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4611 : Intrinsic::x86_avx2_psrl_w;
4616 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4617 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4618 : Intrinsic::x86_avx512_psrl_d_512;
4619 else if (
Size ==
'q')
4620 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4621 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4622 : Intrinsic::x86_avx512_psrl_q_512;
4623 else if (
Size ==
'w')
4624 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4625 : Intrinsic::x86_avx512_psrl_w_512;
4631 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4632 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4633 bool IsVariable = Name[16] ==
'v';
4634 char Size = Name[16] ==
'.' ? Name[17]
4635 : Name[17] ==
'.' ? Name[18]
4636 : Name[18] ==
'.' ? Name[19]
4640 if (IsVariable && Name[17] !=
'.') {
4641 if (
Size ==
's' && Name[17] ==
'4')
4642 IID = Intrinsic::x86_avx2_psrav_d;
4643 else if (
Size ==
's' && Name[17] ==
'8')
4644 IID = Intrinsic::x86_avx2_psrav_d_256;
4645 else if (
Size ==
'h' && Name[17] ==
'8')
4646 IID = Intrinsic::x86_avx512_psrav_w_128;
4647 else if (
Size ==
'h' && Name[17] ==
'1')
4648 IID = Intrinsic::x86_avx512_psrav_w_256;
4649 else if (Name[17] ==
'3' && Name[18] ==
'2')
4650 IID = Intrinsic::x86_avx512_psrav_w_512;
4653 }
else if (Name.ends_with(
".128")) {
4655 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4656 : Intrinsic::x86_sse2_psra_d;
4657 else if (
Size ==
'q')
4658 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4659 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4660 : Intrinsic::x86_avx512_psra_q_128;
4661 else if (
Size ==
'w')
4662 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4663 : Intrinsic::x86_sse2_psra_w;
4666 }
else if (Name.ends_with(
".256")) {
4668 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4669 : Intrinsic::x86_avx2_psra_d;
4670 else if (
Size ==
'q')
4671 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4672 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4673 : Intrinsic::x86_avx512_psra_q_256;
4674 else if (
Size ==
'w')
4675 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4676 : Intrinsic::x86_avx2_psra_w;
4681 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4682 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4683 : Intrinsic::x86_avx512_psra_d_512;
4684 else if (
Size ==
'q')
4685 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4686 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4687 : Intrinsic::x86_avx512_psra_q_512;
4688 else if (
Size ==
'w')
4689 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4690 : Intrinsic::x86_avx512_psra_w_512;
4696 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4698 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4700 }
else if (Name.ends_with(
".movntdqa")) {
4704 LoadInst *LI = Builder.CreateAlignedLoad(
4709 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4710 Name.starts_with(
"fma.vfmsub.") ||
4711 Name.starts_with(
"fma.vfnmadd.") ||
4712 Name.starts_with(
"fma.vfnmsub.")) {
4713 bool NegMul = Name[6] ==
'n';
4714 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4715 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4726 if (NegMul && !IsScalar)
4727 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4728 if (NegMul && IsScalar)
4729 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4731 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4733 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4737 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4745 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4749 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4750 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4751 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4752 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4753 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4754 bool IsMask3 = Name[11] ==
'3';
4755 bool IsMaskZ = Name[11] ==
'z';
4757 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4758 bool NegMul = Name[2] ==
'n';
4759 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4765 if (NegMul && (IsMask3 || IsMaskZ))
4766 A = Builder.CreateFNeg(
A);
4767 if (NegMul && !(IsMask3 || IsMaskZ))
4768 B = Builder.CreateFNeg(
B);
4770 C = Builder.CreateFNeg(
C);
4772 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4773 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4774 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4781 if (Name.back() ==
'd')
4782 IID = Intrinsic::x86_avx512_vfmadd_f64;
4784 IID = Intrinsic::x86_avx512_vfmadd_f32;
4785 Rep = Builder.CreateIntrinsic(IID,
Ops);
4787 Rep = Builder.CreateFMA(
A,
B,
C);
4796 if (NegAcc && IsMask3)
4801 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4803 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4804 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4805 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4806 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4807 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4808 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4809 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4810 bool IsMask3 = Name[11] ==
'3';
4811 bool IsMaskZ = Name[11] ==
'z';
4813 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4814 bool NegMul = Name[2] ==
'n';
4815 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4821 if (NegMul && (IsMask3 || IsMaskZ))
4822 A = Builder.CreateFNeg(
A);
4823 if (NegMul && !(IsMask3 || IsMaskZ))
4824 B = Builder.CreateFNeg(
B);
4826 C = Builder.CreateFNeg(
C);
4833 if (Name[Name.size() - 5] ==
's')
4834 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4836 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4840 Rep = Builder.CreateFMA(
A,
B,
C);
4848 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4852 if (VecWidth == 128 && EltWidth == 32)
4853 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4854 else if (VecWidth == 256 && EltWidth == 32)
4855 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4856 else if (VecWidth == 128 && EltWidth == 64)
4857 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4858 else if (VecWidth == 256 && EltWidth == 64)
4859 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4865 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4866 Rep = Builder.CreateIntrinsic(IID,
Ops);
4867 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4868 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4869 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4870 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4871 bool IsMask3 = Name[11] ==
'3';
4872 bool IsMaskZ = Name[11] ==
'z';
4874 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4875 bool IsSubAdd = Name[3] ==
's';
4879 if (Name[Name.size() - 5] ==
's')
4880 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4882 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4887 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4889 Rep = Builder.CreateIntrinsic(IID,
Ops);
4898 Value *Odd = Builder.CreateCall(FMA,
Ops);
4899 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4900 Value *Even = Builder.CreateCall(FMA,
Ops);
4906 for (
int i = 0; i != NumElts; ++i)
4907 Idxs[i] = i + (i % 2) * NumElts;
4909 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4917 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4918 Name.starts_with(
"avx512.maskz.pternlog.")) {
4919 bool ZeroMask = Name[11] ==
'z';
4923 if (VecWidth == 128 && EltWidth == 32)
4924 IID = Intrinsic::x86_avx512_pternlog_d_128;
4925 else if (VecWidth == 256 && EltWidth == 32)
4926 IID = Intrinsic::x86_avx512_pternlog_d_256;
4927 else if (VecWidth == 512 && EltWidth == 32)
4928 IID = Intrinsic::x86_avx512_pternlog_d_512;
4929 else if (VecWidth == 128 && EltWidth == 64)
4930 IID = Intrinsic::x86_avx512_pternlog_q_128;
4931 else if (VecWidth == 256 && EltWidth == 64)
4932 IID = Intrinsic::x86_avx512_pternlog_q_256;
4933 else if (VecWidth == 512 && EltWidth == 64)
4934 IID = Intrinsic::x86_avx512_pternlog_q_512;
4940 Rep = Builder.CreateIntrinsic(IID, Args);
4944 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4945 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4946 bool ZeroMask = Name[11] ==
'z';
4947 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4950 if (VecWidth == 128 && !
High)
4951 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4952 else if (VecWidth == 256 && !
High)
4953 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4954 else if (VecWidth == 512 && !
High)
4955 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4956 else if (VecWidth == 128 &&
High)
4957 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4958 else if (VecWidth == 256 &&
High)
4959 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4960 else if (VecWidth == 512 &&
High)
4961 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4967 Rep = Builder.CreateIntrinsic(IID, Args);
4971 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4972 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4973 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4974 bool ZeroMask = Name[11] ==
'z';
4975 bool IndexForm = Name[17] ==
'i';
4977 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4978 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4979 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4980 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4981 bool ZeroMask = Name[11] ==
'z';
4982 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4985 if (VecWidth == 128 && !IsSaturating)
4986 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4987 else if (VecWidth == 256 && !IsSaturating)
4988 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4989 else if (VecWidth == 512 && !IsSaturating)
4990 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4991 else if (VecWidth == 128 && IsSaturating)
4992 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4993 else if (VecWidth == 256 && IsSaturating)
4994 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4995 else if (VecWidth == 512 && IsSaturating)
4996 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5006 if (Args[1]->
getType()->isVectorTy() &&
5009 ->isIntegerTy(32) &&
5010 Args[2]->
getType()->isVectorTy() &&
5013 ->isIntegerTy(32)) {
5014 Type *NewArgType =
nullptr;
5015 if (VecWidth == 128)
5017 else if (VecWidth == 256)
5019 else if (VecWidth == 512)
5025 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5026 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5029 Rep = Builder.CreateIntrinsic(IID, Args);
5033 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5034 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5035 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5036 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5037 bool ZeroMask = Name[11] ==
'z';
5038 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5041 if (VecWidth == 128 && !IsSaturating)
5042 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5043 else if (VecWidth == 256 && !IsSaturating)
5044 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5045 else if (VecWidth == 512 && !IsSaturating)
5046 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5047 else if (VecWidth == 128 && IsSaturating)
5048 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5049 else if (VecWidth == 256 && IsSaturating)
5050 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5051 else if (VecWidth == 512 && IsSaturating)
5052 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5062 if (Args[1]->
getType()->isVectorTy() &&
5065 ->isIntegerTy(32) &&
5066 Args[2]->
getType()->isVectorTy() &&
5069 ->isIntegerTy(32)) {
5070 Type *NewArgType =
nullptr;
5071 if (VecWidth == 128)
5073 else if (VecWidth == 256)
5075 else if (VecWidth == 512)
5081 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5082 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5085 Rep = Builder.CreateIntrinsic(IID, Args);
5089 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5090 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5091 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5093 if (Name[0] ==
'a' && Name.back() ==
'2')
5094 IID = Intrinsic::x86_addcarry_32;
5095 else if (Name[0] ==
'a' && Name.back() ==
'4')
5096 IID = Intrinsic::x86_addcarry_64;
5097 else if (Name[0] ==
's' && Name.back() ==
'2')
5098 IID = Intrinsic::x86_subborrow_32;
5099 else if (Name[0] ==
's' && Name.back() ==
'4')
5100 IID = Intrinsic::x86_subborrow_64;
5107 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5110 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5113 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5117 }
else if (Name.starts_with(
"avx512.mask.") &&
5120 }
else if (Name.starts_with(
"bmi.pdep.")) {
5122 }
else if (Name.starts_with(
"bmi.pext.")) {
5132 if (Name.starts_with(
"neon.bfcvt")) {
5133 if (Name.starts_with(
"neon.bfcvtn2")) {
5135 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5137 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5138 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5141 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5142 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5144 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5148 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5149 return Builder.CreateShuffleVector(
5152 return Builder.CreateFPTrunc(CI->
getOperand(0),
5155 }
else if (Name.starts_with(
"sve.fcvt")) {
5158 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5159 .
Case(
"sve.fcvtnt.bf16f32",
5160 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5172 if (Args[1]->
getType() != BadPredTy)
5175 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5176 BadPredTy, Args[1]);
5177 Args[1] = Builder.CreateIntrinsic(
5178 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5180 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5184 if (Name ==
"neon.vcvtfp2hf")
5185 return Builder.CreateBitCast(
5186 Builder.CreateFPTrunc(
5190 if (Name ==
"neon.vcvthf2fp")
5191 return Builder.CreateFPExt(
5192 Builder.CreateBitCast(
5202 if (Name ==
"mve.vctp64.old") {
5205 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5208 Value *C1 = Builder.CreateIntrinsic(
5209 Intrinsic::arm_mve_pred_v2i,
5211 return Builder.CreateIntrinsic(
5212 Intrinsic::arm_mve_pred_i2v,
5214 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5215 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5216 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5217 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5219 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5220 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5221 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5222 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5224 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5225 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5226 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5227 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5228 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5229 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5230 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5231 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5232 std::vector<Type *> Tys;
5236 case Intrinsic::arm_mve_mull_int_predicated:
5237 case Intrinsic::arm_mve_vqdmull_predicated:
5238 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5241 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5242 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5243 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5247 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5251 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5255 case Intrinsic::arm_cde_vcx1q_predicated:
5256 case Intrinsic::arm_cde_vcx1qa_predicated:
5257 case Intrinsic::arm_cde_vcx2q_predicated:
5258 case Intrinsic::arm_cde_vcx2qa_predicated:
5259 case Intrinsic::arm_cde_vcx3q_predicated:
5260 case Intrinsic::arm_cde_vcx3qa_predicated:
5267 std::vector<Value *>
Ops;
5269 Type *Ty =
Op->getType();
5270 if (Ty->getScalarSizeInBits() == 1) {
5271 Value *C1 = Builder.CreateIntrinsic(
5272 Intrinsic::arm_mve_pred_v2i,
5274 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5279 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5294 auto UpgradeLegacyWMMAIUIntrinsicCall =
5299 Args.push_back(Builder.getFalse());
5303 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5310 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5314 NewCall->copyMetadata(*CI);
5318 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5319 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5320 "intrinsic should have 7 arguments");
5323 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5325 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5326 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5327 "intrinsic should have 8 arguments");
5332 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5335 switch (
F->getIntrinsicID()) {
5338 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5339 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5340 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5341 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5342 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5343 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5358 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5361 F->getParent(),
F->getIntrinsicID(), Overloads);
5366 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5370 NewCall->copyMetadata(*CI);
5371 NewCall->takeName(CI);
5376 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5382 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5383 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5391 if (Name.starts_with(
"addrspacecast.nonnull")) {
5394 Value *ASC = Builder.CreateAddrSpaceCast(
5417 if (NumOperands < 3)
5430 bool IsVolatile =
false;
5434 if (NumOperands > 3)
5439 if (NumOperands > 5) {
5441 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5455 if (VT->getElementType()->isIntegerTy(16)) {
5458 Val = Builder.CreateBitCast(Val, AsBF16);
5466 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5468 unsigned AddrSpace = PtrTy->getAddressSpace();
5471 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5473 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5478 MDNode *RangeNotPrivate =
5481 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5487 return Builder.CreateBitCast(RMW, RetTy);
5508 return MAV->getMetadata();
5517 if (Name ==
"label") {
5519 }
else if (Name ==
"assign") {
5526 }
else if (Name ==
"declare") {
5530 }
else if (Name ==
"addr") {
5540 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5541 }
else if (Name ==
"value") {
5544 unsigned ExprOp = 2;
5559 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5567 int64_t OffsetVal =
Offset->getSExtValue();
5568 return Builder.CreateIntrinsic(OffsetVal >= 0
5569 ? Intrinsic::vector_splice_left
5570 : Intrinsic::vector_splice_right,
5572 {CI->getArgOperand(0), CI->getArgOperand(1),
5573 Builder.getInt32(std::abs(OffsetVal))});
5578 if (Name.starts_with(
"to.fp16")) {
5580 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5581 return Builder.CreateBitCast(Cast, CI->
getType());
5584 if (Name.starts_with(
"from.fp16")) {
5586 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5587 return Builder.CreateFPExt(Cast, CI->
getType());
5646 else if (Opcode == Instruction::ICmp)
5649 else if (Opcode == Instruction::FCmp)
5652 else if (Opcode == Instruction::Select)
5657 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5669 if (Defaults.empty())
5672 unsigned OldArgCount = CI->
arg_size();
5673 unsigned NewArgCount = NewFn->
arg_size();
5675 if (OldArgCount < FirstDefault)
5679 if (OldArgCount > NewArgCount)
5684 if (OldArgCount == NewArgCount) {
5696 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5697 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5698 "missing argument outside the default range");
5699 Type *ParamTy = NewFT->getParamType(Idx);
5704 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5710 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5742 if (!Name.consume_front(
"llvm."))
5745 bool IsX86 = Name.consume_front(
"x86.");
5746 bool IsNVVM = Name.consume_front(
"nvvm.");
5747 bool IsAArch64 = Name.consume_front(
"aarch64.");
5748 bool IsARM = Name.consume_front(
"arm.");
5749 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5750 bool IsDbg = Name.consume_front(
"dbg.");
5752 (Name.consume_front(
"experimental.vector.splice") ||
5753 Name.consume_front(
"vector.splice")) &&
5754 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5755 Value *Rep =
nullptr;
5757 if (!IsX86 && Name ==
"stackprotectorcheck") {
5759 }
else if (IsNVVM) {
5763 }
else if (IsAArch64) {
5767 }
else if (IsAMDGCN) {
5771 }
else if (IsOldSplice) {
5773 }
else if (Name.consume_front(
"convert.")) {
5775 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5790 const auto &DefaultCase = [&]() ->
void {
5798 "Unknown function for CallBase upgrade and isn't just a name change");
5806 "Return type must have changed");
5807 assert(OldST->getNumElements() ==
5809 "Must have same number of elements");
5812 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5815 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5816 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5817 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5838 case Intrinsic::arm_neon_vst1:
5839 case Intrinsic::arm_neon_vst2:
5840 case Intrinsic::arm_neon_vst3:
5841 case Intrinsic::arm_neon_vst4:
5842 case Intrinsic::arm_neon_vst2lane:
5843 case Intrinsic::arm_neon_vst3lane:
5844 case Intrinsic::arm_neon_vst4lane: {
5846 NewCall = Builder.CreateCall(NewFn, Args);
5849 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5850 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5851 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5856 NewCall = Builder.CreateCall(NewFn, Args);
5859 case Intrinsic::aarch64_sve_ld3_sret:
5860 case Intrinsic::aarch64_sve_ld4_sret:
5861 case Intrinsic::aarch64_sve_ld2_sret: {
5869 Name = Name.substr(5);
5876 unsigned MinElts = RetTy->getMinNumElements() /
N;
5878 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5880 for (
unsigned I = 0;
I <
N;
I++) {
5881 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5882 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5888 case Intrinsic::coro_end_async:
5889 case Intrinsic::coro_end: {
5891 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5893 NewCall = Builder.CreateCall(NewFn, Args);
5898 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5899 Value *InRamp = Builder.CreateCall(IsInRamp);
5909 case Intrinsic::vector_extract: {
5911 Name = Name.substr(5);
5912 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5917 unsigned MinElts = RetTy->getMinNumElements();
5920 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5924 case Intrinsic::vector_insert: {
5926 Name = Name.substr(5);
5927 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5931 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5936 NewCall = Builder.CreateCall(
5940 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5946 assert(
N > 1 &&
"Create is expected to be between 2-4");
5949 unsigned MinElts = RetTy->getMinNumElements() /
N;
5950 for (
unsigned I = 0;
I <
N;
I++) {
5952 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5959 case Intrinsic::arm_neon_bfdot:
5960 case Intrinsic::arm_neon_bfmmla:
5961 case Intrinsic::arm_neon_bfmlalb:
5962 case Intrinsic::arm_neon_bfmlalt:
5963 case Intrinsic::aarch64_neon_bfdot:
5964 case Intrinsic::aarch64_neon_bfmmla:
5965 case Intrinsic::aarch64_neon_bfmlalb:
5966 case Intrinsic::aarch64_neon_bfmlalt: {
5969 "Mismatch between function args and call args");
5970 size_t OperandWidth =
5972 assert((OperandWidth == 64 || OperandWidth == 128) &&
5973 "Unexpected operand width");
5975 auto Iter = CI->
args().begin();
5976 Args.push_back(*Iter++);
5977 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5978 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5979 NewCall = Builder.CreateCall(NewFn, Args);
5983 case Intrinsic::bitreverse:
5984 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5987 case Intrinsic::ctlz:
5988 case Intrinsic::cttz: {
5995 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5999 case Intrinsic::objectsize: {
6000 Value *NullIsUnknownSize =
6004 NewCall = Builder.CreateCall(
6009 case Intrinsic::ctpop:
6010 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6012 case Intrinsic::dbg_value: {
6014 Name = Name.substr(5);
6016 if (Name.starts_with(
"dbg.addr")) {
6030 if (
Offset->isNullValue()) {
6031 NewCall = Builder.CreateCall(
6040 case Intrinsic::ptr_annotation:
6048 NewCall = Builder.CreateCall(
6057 case Intrinsic::var_annotation:
6064 NewCall = Builder.CreateCall(
6073 case Intrinsic::riscv_aes32dsi:
6074 case Intrinsic::riscv_aes32dsmi:
6075 case Intrinsic::riscv_aes32esi:
6076 case Intrinsic::riscv_aes32esmi:
6077 case Intrinsic::riscv_sm4ks:
6078 case Intrinsic::riscv_sm4ed: {
6088 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6089 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6095 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6096 Value *Res = NewCall;
6098 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6104 case Intrinsic::nvvm_mapa_shared_cluster: {
6108 Value *Res = NewCall;
6109 Res = Builder.CreateAddrSpaceCast(
6116 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6118 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6120 Args[0] = Builder.CreateAddrSpaceCast(
6124 Args.push_back(Builder.getInt32(0));
6126 NewCall = Builder.CreateCall(NewFn, Args);
6132 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6137 for (
unsigned I = 0;
I < 4; ++
I)
6139 Args.push_back(Builder.getInt32(0));
6140 Args.push_back(Builder.getInt32(0));
6143 Args.push_back(Builder.getInt1(
false));
6144 Args.push_back(Builder.getInt32(0));
6146 NewCall = Builder.CreateCall(NewFn, Args);
6152 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6155 Args[0] = Builder.CreateAddrSpaceCast(
6158 NewCall = Builder.CreateCall(NewFn, Args);
6165#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6166 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6168#undef G2S_CLUSTER_CASE
6173 Args[0] = Builder.CreateAddrSpaceCast(
6179 Args.push_back(Builder.getInt32(0));
6181 NewCall = Builder.CreateCall(NewFn, Args);
6188#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6189 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6197 "expected only the trailing flag_valid_pattern to be missing");
6198 Args.push_back(Builder.getInt32(0));
6200 NewCall = Builder.CreateCall(NewFn, Args);
6206#undef NVVM_TMA_G2S_MODES
6209 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6210 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6211 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6212 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6213 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6214 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6215 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6216 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6218 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6222 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6223 NewCall = Builder.CreateCall(NewFn, Args);
6226 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6227 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6228 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6229 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6232 Builder.getFalse()});
6234 case Intrinsic::nvvm_mbarrier_init: {
6238 if (Args.size() == 2)
6239 Args.push_back(Builder.getInt32(0));
6240 NewCall = Builder.CreateCall(NewFn, Args);
6243 case Intrinsic::riscv_sha256sig0:
6244 case Intrinsic::riscv_sha256sig1:
6245 case Intrinsic::riscv_sha256sum0:
6246 case Intrinsic::riscv_sha256sum1:
6247 case Intrinsic::riscv_sm3p0:
6248 case Intrinsic::riscv_sm3p1: {
6255 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6257 NewCall = Builder.CreateCall(NewFn, Arg);
6259 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6266 case Intrinsic::x86_xop_vfrcz_ss:
6267 case Intrinsic::x86_xop_vfrcz_sd:
6268 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6271 case Intrinsic::x86_xop_vpermil2pd:
6272 case Intrinsic::x86_xop_vpermil2ps:
6273 case Intrinsic::x86_xop_vpermil2pd_256:
6274 case Intrinsic::x86_xop_vpermil2ps_256: {
6278 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6279 NewCall = Builder.CreateCall(NewFn, Args);
6283 case Intrinsic::x86_sse41_ptestc:
6284 case Intrinsic::x86_sse41_ptestz:
6285 case Intrinsic::x86_sse41_ptestnzc: {
6299 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6300 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6302 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6306 case Intrinsic::x86_rdtscp: {
6312 NewCall = Builder.CreateCall(NewFn);
6314 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6317 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6325 case Intrinsic::x86_sse41_insertps:
6326 case Intrinsic::x86_sse41_dppd:
6327 case Intrinsic::x86_sse41_dpps:
6328 case Intrinsic::x86_sse41_mpsadbw:
6329 case Intrinsic::x86_avx_dp_ps_256:
6330 case Intrinsic::x86_avx2_mpsadbw: {
6336 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6337 NewCall = Builder.CreateCall(NewFn, Args);
6341 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6342 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6343 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6344 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6345 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6346 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6352 NewCall = Builder.CreateCall(NewFn, Args);
6361 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6362 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6363 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6364 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6365 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6366 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6370 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6371 Args[1] = Builder.CreateBitCast(
6374 NewCall = Builder.CreateCall(NewFn, Args);
6375 Value *Res = Builder.CreateBitCast(
6383 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6384 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6385 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6389 Args[1] = Builder.CreateBitCast(
6391 Args[2] = Builder.CreateBitCast(
6394 NewCall = Builder.CreateCall(NewFn, Args);
6398 case Intrinsic::thread_pointer: {
6399 NewCall = Builder.CreateCall(NewFn, {});
6403 case Intrinsic::memcpy:
6404 case Intrinsic::memmove:
6405 case Intrinsic::memset: {
6421 NewCall = Builder.CreateCall(NewFn, Args);
6424 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6425 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6426 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6431 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6434 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6438 case Intrinsic::masked_load:
6439 case Intrinsic::masked_gather:
6440 case Intrinsic::masked_store:
6441 case Intrinsic::masked_scatter: {
6447 auto GetMaybeAlign = [](
Value *
Op) {
6449 uint64_t Val = CI->getZExtValue();
6457 auto GetAlign = [&](
Value *
Op) {
6466 case Intrinsic::masked_load:
6467 NewCall = Builder.CreateMaskedLoad(
6471 case Intrinsic::masked_gather:
6472 NewCall = Builder.CreateMaskedGather(
6478 case Intrinsic::masked_store:
6479 NewCall = Builder.CreateMaskedStore(
6483 case Intrinsic::masked_scatter:
6484 NewCall = Builder.CreateMaskedScatter(
6486 DL.getValueOrABITypeAlignment(
6500 case Intrinsic::lifetime_start:
6501 case Intrinsic::lifetime_end: {
6513 NewCall = Builder.CreateLifetimeStart(Ptr);
6515 NewCall = Builder.CreateLifetimeEnd(Ptr);
6524 case Intrinsic::x86_avx512_vpdpbusd_128:
6525 case Intrinsic::x86_avx512_vpdpbusd_256:
6526 case Intrinsic::x86_avx512_vpdpbusd_512:
6527 case Intrinsic::x86_avx512_vpdpbusds_128:
6528 case Intrinsic::x86_avx512_vpdpbusds_256:
6529 case Intrinsic::x86_avx512_vpdpbusds_512:
6530 case Intrinsic::x86_avx2_vpdpbssd_128:
6531 case Intrinsic::x86_avx2_vpdpbssd_256:
6532 case Intrinsic::x86_avx10_vpdpbssd_512:
6533 case Intrinsic::x86_avx2_vpdpbssds_128:
6534 case Intrinsic::x86_avx2_vpdpbssds_256:
6535 case Intrinsic::x86_avx10_vpdpbssds_512:
6536 case Intrinsic::x86_avx2_vpdpbsud_128:
6537 case Intrinsic::x86_avx2_vpdpbsud_256:
6538 case Intrinsic::x86_avx10_vpdpbsud_512:
6539 case Intrinsic::x86_avx2_vpdpbsuds_128:
6540 case Intrinsic::x86_avx2_vpdpbsuds_256:
6541 case Intrinsic::x86_avx10_vpdpbsuds_512:
6542 case Intrinsic::x86_avx2_vpdpbuud_128:
6543 case Intrinsic::x86_avx2_vpdpbuud_256:
6544 case Intrinsic::x86_avx10_vpdpbuud_512:
6545 case Intrinsic::x86_avx2_vpdpbuuds_128:
6546 case Intrinsic::x86_avx2_vpdpbuuds_256:
6547 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6552 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6553 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6555 NewCall = Builder.CreateCall(NewFn, Args);
6558 case Intrinsic::x86_avx512_vpdpwssd_128:
6559 case Intrinsic::x86_avx512_vpdpwssd_256:
6560 case Intrinsic::x86_avx512_vpdpwssd_512:
6561 case Intrinsic::x86_avx512_vpdpwssds_128:
6562 case Intrinsic::x86_avx512_vpdpwssds_256:
6563 case Intrinsic::x86_avx512_vpdpwssds_512:
6564 case Intrinsic::x86_avx2_vpdpwsud_128:
6565 case Intrinsic::x86_avx2_vpdpwsud_256:
6566 case Intrinsic::x86_avx10_vpdpwsud_512:
6567 case Intrinsic::x86_avx2_vpdpwsuds_128:
6568 case Intrinsic::x86_avx2_vpdpwsuds_256:
6569 case Intrinsic::x86_avx10_vpdpwsuds_512:
6570 case Intrinsic::x86_avx2_vpdpwusd_128:
6571 case Intrinsic::x86_avx2_vpdpwusd_256:
6572 case Intrinsic::x86_avx10_vpdpwusd_512:
6573 case Intrinsic::x86_avx2_vpdpwusds_128:
6574 case Intrinsic::x86_avx2_vpdpwusds_256:
6575 case Intrinsic::x86_avx10_vpdpwusds_512:
6576 case Intrinsic::x86_avx2_vpdpwuud_128:
6577 case Intrinsic::x86_avx2_vpdpwuud_256:
6578 case Intrinsic::x86_avx10_vpdpwuud_512:
6579 case Intrinsic::x86_avx2_vpdpwuuds_128:
6580 case Intrinsic::x86_avx2_vpdpwuuds_256:
6581 case Intrinsic::x86_avx10_vpdpwuuds_512:
6586 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6587 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6589 NewCall = Builder.CreateCall(NewFn, Args);
6592 assert(NewCall &&
"Should have either set this variable or returned through "
6593 "the default case");
6600 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6614 F->eraseFromParent();
6620 if (NumOperands == 0)
6628 if (NumOperands == 3) {
6632 Metadata *Elts2[] = {ScalarType, ScalarType,
6646 if (
Opc != Instruction::BitCast)
6650 Type *SrcTy = V->getType();
6667 if (
Opc != Instruction::BitCast)
6670 Type *SrcTy =
C->getType();
6687 if (Flag.getNumOperands() < 3)
6688 return std::nullopt;
6690 return Name->getString();
6691 return std::nullopt;
6705 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6706 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6707 if (auto Name = getModuleFlagNameSafely(*Flag))
6708 return *Name ==
"Debug Info Version";
6711 if (OpIt != ModFlags->op_end()) {
6712 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6719 bool BrokenDebugInfo =
false;
6722 if (!BrokenDebugInfo)
6728 M.getContext().diagnose(Diag);
6735 M.getContext().diagnose(DiagVersion);
6745 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6748 if (
F->hasFnAttribute(Attr)) {
6751 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6753 auto [Part, Rest] = S.
split(
',');
6759 const unsigned Dim = DimC -
'x';
6760 assert(Dim < 3 &&
"Unexpected dim char");
6770 F->addFnAttr(Attr, NewAttr);
6774 return S ==
"x" || S ==
"y" || S ==
"z";
6779 if (
K ==
"kernel") {
6791 const unsigned Idx = (AlignIdxValuePair >> 16);
6792 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6797 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6802 if (
K ==
"minctasm") {
6807 if (
K ==
"maxnreg") {
6812 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6816 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6820 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6824 if (
K ==
"grid_constant") {
6839 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6846 if (!SeenNodes.
insert(MD).second)
6853 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6860 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6862 const MDOperand &V = MD->getOperand(j + 1);
6868 if (NewOperands.
size() > 1)
6881 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6882 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6883 if (ModRetainReleaseMarker) {
6889 ID->getString().split(ValueComp,
"#");
6890 if (ValueComp.
size() == 2) {
6891 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6895 M.eraseNamedMetadata(ModRetainReleaseMarker);
6906 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6932 bool InvalidCast =
false;
6934 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6947 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6949 Args.push_back(Arg);
6956 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6961 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6974 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6982 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6983 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6984 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6985 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6986 {
"objc_autoreleaseReturnValue",
6987 llvm::Intrinsic::objc_autoreleaseReturnValue},
6988 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6989 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6990 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6991 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6992 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6993 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6994 {
"objc_release", llvm::Intrinsic::objc_release},
6995 {
"objc_retain", llvm::Intrinsic::objc_retain},
6996 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6997 {
"objc_retainAutoreleaseReturnValue",
6998 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6999 {
"objc_retainAutoreleasedReturnValue",
7000 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7001 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7002 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7003 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7004 {
"objc_unsafeClaimAutoreleasedReturnValue",
7005 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7006 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7007 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7008 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7009 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7010 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7011 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7012 {
"objc_arc_annotation_topdown_bbstart",
7013 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7014 {
"objc_arc_annotation_topdown_bbend",
7015 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7016 {
"objc_arc_annotation_bottomup_bbstart",
7017 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7018 {
"objc_arc_annotation_bottomup_bbend",
7019 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7021 for (
auto &
I : RuntimeFuncs)
7022 UpgradeToIntrinsic(
I.first,
I.second);
7046 std::optional<bool> UseAddressDisc;
7049 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7050 for (
const MDNode *Flag : ModFlags->operands()) {
7052 if (Name && (*Name ==
"ptrauth-init-fini" ||
7053 *Name ==
"ptrauth-init-fini-address-discrimination"))
7058 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7059 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7060 constexpr unsigned ExpectedAddressMarker = 1;
7063 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7066 bool HasAddressDisc;
7067 if (!CPA->hasAddressDiscriminator())
7068 HasAddressDisc =
false;
7069 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7070 HasAddressDisc =
true;
7074 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7077 UseAddressDisc = HasAddressDisc;
7078 return CPA->getPointer();
7082 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7085 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7087 if (!GV || !GV->hasInitializer())
7091 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7094 std::vector<Constant *> NewStructors;
7095 NewStructors.reserve(OldStructorsArray->getNumOperands());
7097 for (
Use &U : OldStructorsArray->operands()) {
7106 Func = UpgradeSinglePointer(Func);
7110 NewStructors.push_back(
7119 if (GlobalArraysToUpgrade.
empty())
7121 assert(UseAddressDisc.has_value());
7123 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7124 GV->setInitializer(NewInit);
7127 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7137 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7141 bool HasObjCFlag =
false, HasClassProperties =
false;
7142 bool HasSwiftVersionFlag =
false;
7143 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7150 if (
Op->getNumOperands() != 3)
7164 if (ID->getString() ==
"Objective-C Image Info Version")
7166 if (ID->getString() ==
"Objective-C Class Properties")
7167 HasClassProperties =
true;
7169 if (ID->getString() ==
"PIC Level") {
7170 if (
auto *Behavior =
7172 uint64_t V = Behavior->getLimitedValue();
7178 if (ID->getString() ==
"PIE Level")
7179 if (
auto *Behavior =
7186 if (ID->getString() ==
"branch-target-enforcement" ||
7187 ID->getString().starts_with(
"sign-return-address")) {
7188 if (
auto *Behavior =
7194 Op->getOperand(1),
Op->getOperand(2)};
7204 if (ID->getString() ==
"Objective-C Image Info Section") {
7207 Value->getString().split(ValueComp,
" ");
7208 if (ValueComp.
size() != 1) {
7209 std::string NewValue;
7210 for (
auto &S : ValueComp)
7211 NewValue += S.str();
7222 if (ID->getString() ==
"Objective-C Garbage Collection") {
7225 assert(Md->getValue() &&
"Expected non-empty metadata");
7226 auto Type = Md->getValue()->getType();
7229 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7230 if ((Val & 0xff) != Val) {
7231 HasSwiftVersionFlag =
true;
7232 SwiftABIVersion = (Val & 0xff00) >> 8;
7233 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7234 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7245 if (ID->getString() ==
"amdgpu_code_object_version") {
7248 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7257 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7286 if (HasObjCFlag && !HasClassProperties) {
7292 if (HasSwiftVersionFlag) {
7296 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7298 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7306 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7310 auto MatchesVersion = [](
const MDNode *
Op) {
7311 return Op->getNumOperands() >= 3 &&
7325 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7326 assert(
Op->getNumOperands() >= 2 &&
7327 "Expected at least 2 operands - name and linkage type");
7339 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7350 auto TrimSpaces = [](
StringRef Section) -> std::string {
7352 Section.split(Components,
',');
7357 for (
auto Component : Components)
7358 OS <<
',' << Component.trim();
7363 for (
auto &GV : M.globals()) {
7364 if (!GV.hasSection())
7369 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7374 GV.setSection(TrimSpaces(Section));
7390struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7391 StrictFPUpgradeVisitor() =
default;
7394 if (!
Call.isStrictFP())
7400 Call.removeFnAttr(Attribute::StrictFP);
7401 Call.addFnAttr(Attribute::NoBuiltin);
7406struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7407 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7408 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7410 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7425 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7426 StrictFPUpgradeVisitor SFPV;
7431 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7432 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7433 for (
auto &Arg :
F.args())
7435 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7437 bool AddingAttrs =
false, RemovingAttrs =
false;
7438 AttrBuilder AttrsToAdd(
F.getContext());
7443 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7444 A.isValid() &&
A.isStringAttribute()) {
7445 F.setSection(
A.getValueAsString());
7447 RemovingAttrs =
true;
7451 A.isValid() &&
A.isStringAttribute()) {
7454 AddingAttrs = RemovingAttrs =
true;
7457 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7458 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7460 RemovingAttrs =
true;
7461 if (
A.getValueAsString() ==
"true") {
7462 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7471 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7474 if (
A.getValueAsBool()) {
7475 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7481 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7482 RemovingAttrs =
true;
7489 bool HandleDenormalMode =
false;
7491 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7494 DenormalFPMath = ParsedMode;
7496 AddingAttrs = RemovingAttrs =
true;
7497 HandleDenormalMode =
true;
7501 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7505 DenormalFPMathF32 = ParsedMode;
7507 AddingAttrs = RemovingAttrs =
true;
7508 HandleDenormalMode =
true;
7512 if (HandleDenormalMode)
7513 AttrsToAdd.addDenormalFPEnvAttr(
7517 F.removeFnAttrs(AttrsToRemove);
7520 F.addFnAttrs(AttrsToAdd);
7526 if (!
F.hasFnAttribute(FnAttrName))
7527 F.addFnAttr(FnAttrName,
Value);
7534 if (!
F.hasFnAttribute(FnAttrName)) {
7536 F.addFnAttr(FnAttrName);
7538 auto A =
F.getFnAttribute(FnAttrName);
7539 if (
"false" ==
A.getValueAsString())
7540 F.removeFnAttr(FnAttrName);
7541 else if (
"true" ==
A.getValueAsString()) {
7542 F.removeFnAttr(FnAttrName);
7543 F.addFnAttr(FnAttrName);
7549 Triple T(M.getTargetTriple());
7550 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7553 uint64_t BTEValue = 0;
7554 uint64_t BPPLRValue = 0;
7555 uint64_t GCSValue = 0;
7556 uint64_t SRAValue = 0;
7557 uint64_t SRAALLValue = 0;
7558 uint64_t SRABKeyValue = 0;
7560 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7564 if (
Op->getNumOperands() != 3)
7573 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7574 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7575 : IDStr ==
"guarded-control-stack" ? &GCSValue
7576 : IDStr ==
"sign-return-address" ? &SRAValue
7577 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7578 : IDStr ==
"sign-return-address-with-bkey"
7584 *ValPtr = CI->getZExtValue();
7590 bool BTE = BTEValue == 1;
7591 bool BPPLR = BPPLRValue == 1;
7592 bool GCS = GCSValue == 1;
7593 bool SRA = SRAValue == 1;
7596 if (SRA && SRAALLValue == 1)
7597 SignTypeValue =
"all";
7600 if (SRA && SRABKeyValue == 1)
7601 SignKeyValue =
"b_key";
7603 for (
Function &
F : M.getFunctionList()) {
7604 if (
F.isDeclaration())
7611 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7612 A.isValid() &&
"none" ==
A.getValueAsString()) {
7613 F.removeFnAttr(
"sign-return-address");
7614 F.removeFnAttr(
"sign-return-address-key");
7630 if (SRAALLValue == 1)
7632 if (SRABKeyValue == 1)
7659 if (
T->getNumOperands() < 1)
7664 if (S->getString().starts_with(
"llvm.vectorizer."))
7670 StringRef OldPrefix =
"llvm.vectorizer.";
7673 if (OldTag ==
"llvm.vectorizer.unroll")
7685 if (
T->getNumOperands() < 1)
7697 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7710 Ops.reserve(
T->getNumOperands());
7711 Ops.push_back(NewTag);
7712 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7713 Ops.push_back(
T->getOperand(
I));
7730 if (
T->isDistinct()) {
7731 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7743 Ops.reserve(
T->getNumOperands());
7754 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7755 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7756 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7759 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7761 auto I =
DL.find(
"-n64-");
7763 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7768 std::string Res =
DL.str();
7771 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7772 Res.append(Res.empty() ?
"G1" :
"-G1");
7780 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7781 Res.append(
"-ni:7:8:9");
7783 if (
DL.ends_with(
"ni:7"))
7785 if (
DL.ends_with(
"ni:7:8"))
7790 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7791 Res.append(
"-p7:160:256:256:32");
7792 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7793 Res.append(
"-p8:128:128:128:48");
7794 constexpr StringRef OldP8(
"-p8:128:128-");
7795 if (
DL.contains(OldP8))
7796 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7797 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7798 Res.append(
"-p9:192:256:256:32");
7803 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7804 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7805 Res.append((
"-" + AS +
":32:32").str());
7810 if (!
DL.contains(
"m:e"))
7811 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7816 if (
T.isSystemZ() && !
DL.empty()) {
7818 if (!
DL.contains(
"-S64"))
7819 return "E-S64" +
DL.drop_front(1).str();
7823 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7826 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7827 if (!
DL.contains(AddrSpaces)) {
7829 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7830 if (R.match(Res, &
Groups))
7836 if (
T.isAArch64()) {
7838 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7839 Res.append(
"-Fn32");
7840 AddPtr32Ptr64AddrSpaces();
7844 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7848 std::string I64 =
"-i64:64";
7849 std::string I128 =
"-i128:128";
7851 size_t Pos = Res.find(I64);
7852 if (Pos !=
size_t(-1))
7853 Res.insert(Pos + I64.size(), I128);
7857 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7858 size_t Pos = Res.find(
"-S128");
7861 Res.insert(Pos,
"-f64:32:64");
7867 AddPtr32Ptr64AddrSpaces();
7875 if (!
T.isOSIAMCU()) {
7876 std::string I128 =
"-i128:128";
7879 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7880 if (R.match(Res, &
Groups))
7888 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7890 auto I =
Ref.find(
"-f80:32-");
7892 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7900 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7903 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7904 B.removeAttribute(
"no-frame-pointer-elim");
7906 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7908 if (FramePointer !=
"all")
7909 FramePointer =
"non-leaf";
7910 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7912 if (!FramePointer.
empty())
7913 B.addAttribute(
"frame-pointer", FramePointer);
7915 A =
B.getAttribute(
"null-pointer-is-valid");
7918 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7919 B.removeAttribute(
"null-pointer-is-valid");
7920 if (NullPointerIsValid)
7921 B.addAttribute(Attribute::NullPointerIsValid);
7924 A =
B.getAttribute(
"uniform-work-group-size");
7928 bool IsTrue = Val ==
"true";
7929 B.removeAttribute(
"uniform-work-group-size");
7931 B.addAttribute(
"uniform-work-group-size");
7942 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static 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)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFAddUpgrade(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 bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static cl::opt< bool > DisableAutoUpgradeDebugInfo("disable-auto-upgrade-debug-info", cl::desc("Disable autoupgrade of debug info"))
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static 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 unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F, StringRef Name)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_CASE(ID_SUFFIX, NAME)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static FCmpInst::Predicate getVPFPPredicateFromMD(const Value *Op)
static GCRegistry::Add< ShadowStackGC > C("shadow-stack", "Very portable GC for uncooperative code generators")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
This file contains constants used for implementing Dwarf debug support.
Module.h This file contains the declarations for the Module class.
const AbstractManglingParser< Derived, Alloc >::OperatorInfo AbstractManglingParser< Derived, Alloc >::Ops[]
static bool isZero(Value *V, const DataLayout &DL, DominatorTree *DT, AssumptionCache *AC)
NVPTX address space definition.
This file contains the definitions of the enumerations and flags associated with NVVM Intrinsics,...
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
static SymbolRef::Type getType(const Symbol *Sym)
LocallyHashedType DenseMapInfo< LocallyHashedType >::Empty
static const X86InstrFMA3Group Groups[]
Class for arbitrary precision integers.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
Class to represent array types.
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Type * getElementType() const
an instruction that atomically reads a memory location, combines it with another value,...
void setVolatile(bool V)
Specify whether this is a volatile RMW or not.
BinOp
This enumeration lists the possible modifications atomicrmw can make.
@ USubCond
Subtract only if no unsigned overflow.
@ Min
*p = old <signed v ? old : v
@ USubSat
*p = usub.sat(old, v) usub.sat matches the behavior of llvm.usub.sat.
@ UIncWrap
Increment one up to a maximum value.
@ Max
*p = old >signed v ? old : v
@ FMin
*p = minnum(old, v) minnum matches the behavior of llvm.minnum.
@ FMax
*p = maxnum(old, v) maxnum matches the behavior of llvm.maxnum.
@ UDecWrap
Decrement one until a minimum value or zero.
bool isFloatingPointOperation() const
This class stores enough information to efficiently remove some attributes from an existing AttrBuild...
AttributeMask & addAttribute(Attribute::AttrKind Val)
Add an attribute to the mask.
Functions, function parameters, and return types can have attributes to indicate how they should be t...
static LLVM_ABI Attribute getWithStackAlignment(LLVMContext &Context, Align Alignment)
static LLVM_ABI Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
void setCallingConv(CallingConv::ID CC)
LLVM_ABI void getOperandBundlesAsDefs(SmallVectorImpl< OperandBundleDef > &Defs) const
Return the list of operand bundles attached to this instruction as a vector of OperandBundleDefs.
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
CallingConv::ID getCallingConv() const
Value * getCalledOperand() const
void setAttributes(AttributeList A)
Set the attributes for this call.
Value * getArgOperand(unsigned i) const
FunctionType * getFunctionType() const
LLVM_ABI Intrinsic::ID getIntrinsicID() const
Returns the intrinsic ID of the intrinsic called or Intrinsic::not_intrinsic if the called function i...
iterator_range< User::op_iterator > args()
Iteration adapter for range-for loops.
void setCalledOperand(Value *V)
unsigned arg_size() const
AttributeList getAttributes() const
Return the attributes for this call.
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
This class represents a function call, abstracting a target machine's calling convention.
void setTailCallKind(TailCallKind TCK)
static LLVM_ABI CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", InsertPosition InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
static LLVM_ABI bool castIsValid(Instruction::CastOps op, Type *SrcTy, Type *DstTy)
This method can be used to determine if a cast from SrcTy to DstTy using Opcode op is valid or not.
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
unsigned getNumParams() const
Return the number of fixed parameters this function type requires.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
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...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
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 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
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
bool is_contained(R &&Range, const E &Element)
Returns true if Element is found in Range.
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
@ 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.