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)
1284 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1287 auto [RedOpName, ShapeName] = Name.split(
'.');
1292 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1293 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1294 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1295 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1296 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1297 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1298 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1299 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1305 if (Name.consume_front(
"mapa.shared.cluster"))
1306 if (
F->getReturnType()->getPointerAddressSpace() ==
1308 return Intrinsic::nvvm_mapa_shared_cluster;
1310 if (Name.consume_front(
"cp.async.bulk.")) {
1313 .
Case(
"global.to.shared.cluster",
1314 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1315 .
Case(
"shared.cta.to.cluster",
1316 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1320 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1330 if (!Name.consume_front(
"tcgen05.commit."))
1333 if (Name.consume_front(
"shared."))
1335 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1336 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1339 if (Name.consume_front(
"mc.shared.")) {
1341 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1345 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1346 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1355 if (
F->arg_size() != 2)
1358 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1359 Name.consume_front(
"tcgen05.alloc."))
1361 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1362 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1365 if (Name.consume_front(
"tcgen05.dealloc."))
1367 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1368 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1375 if (Name.consume_front(
"fma.rn."))
1377 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1378 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1379 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1380 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1383 if (Name.consume_front(
"fmax."))
1385 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1386 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1387 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1388 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1389 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1390 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1391 .
Case(
"ftz.nan.xorsign.abs.bf16",
1392 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1393 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1394 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1395 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1396 .
Case(
"ftz.xorsign.abs.bf16x2",
1397 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1398 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1399 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1400 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1401 .
Case(
"nan.xorsign.abs.bf16x2",
1402 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1403 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1404 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1407 if (Name.consume_front(
"fmin."))
1409 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1410 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1411 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1412 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1413 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1414 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1415 .
Case(
"ftz.nan.xorsign.abs.bf16",
1416 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1417 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1418 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1419 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1420 .
Case(
"ftz.xorsign.abs.bf16x2",
1421 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1422 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1423 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1424 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1425 .
Case(
"nan.xorsign.abs.bf16x2",
1426 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1427 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1428 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1431 if (Name.consume_front(
"neg."))
1433 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1434 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1443 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1448 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1451 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1454 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1455 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1463 if (!Name.consume_front(
"tcgen05.mma."))
1467 if (Name.starts_with(
"ws"))
1470 return F->getIntrinsicID();
1473static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1475 auto [Modifiers,
Type] = Name.rsplit(
'.');
1477 return std::nullopt;
1487 return std::nullopt;
1490 .
Case(
"", Intrinsic::nvvm_fadd)
1491 .
Case(
".ftz", Intrinsic::nvvm_fadd_ftz)
1492 .
Case(
".sat", Intrinsic::nvvm_fadd_sat)
1493 .
Case(
".ftz.sat", Intrinsic::nvvm_fadd_ftz_sat)
1496 return std::nullopt;
1502 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1503 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1504 Name.consume_front(
"param");
1508 if (!Name.consume_front(
"vp."))
1537 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1538 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1545 if (!Name.consume_front(
"vp."))
1565 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1566 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1571 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1583 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1594 if (Name.starts_with(
"to.fp16")) {
1598 FuncTy->getReturnType());
1601 if (Name.starts_with(
"from.fp16")) {
1605 FuncTy->getReturnType());
1617 if (Defaults.empty())
1629 if (
F->arg_size() >= FullDecl->
arg_size())
1634 if (
F->arg_size() < FirstDefault)
1642 bool CanUpgradeDebugIntrinsicsToRecords) {
1643 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1648 if (!Name.consume_front(
"llvm.") || Name.empty())
1654 bool IsArm = Name.consume_front(
"arm.");
1655 if (IsArm || Name.consume_front(
"aarch64.")) {
1661 if (Name.consume_front(
"amdgcn.")) {
1662 if (Name ==
"alignbit") {
1665 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1669 if (Name.consume_front(
"atomic.")) {
1670 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1671 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1680 if (Name.starts_with(
"addrspacecast.nonnull")) {
1687 switch (
F->getIntrinsicID()) {
1691 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1692 if (
F->arg_size() == 7) {
1697 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1698 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1699 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1700 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1701 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1702 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1703 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1704 if (
F->arg_size() == 8) {
1711 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1712 Name.consume_front(
"flat.atomic.")) {
1713 if (Name.starts_with(
"fadd") ||
1715 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1716 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1724 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1729 if (Name.starts_with(
"ldexp.")) {
1732 F->getParent(), Intrinsic::ldexp,
1733 {F->getReturnType(), F->getArg(1)->getType()});
1742 if (
F->arg_size() == 1) {
1743 if (Name.consume_front(
"convert.")) {
1757 F->arg_begin()->getType());
1763 if (Name ==
"coro.end" &&
1764 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1765 CoroEndID = Intrinsic::coro_end;
1766 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1767 CoroEndID = Intrinsic::coro_end_async;
1778 if (Name.consume_front(
"dbg.")) {
1780 if (CanUpgradeDebugIntrinsicsToRecords) {
1781 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1782 Name ==
"declare" || Name ==
"label") {
1791 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1794 Intrinsic::dbg_value);
1801 if (Name.consume_front(
"experimental.vector.")) {
1807 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1808 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1809 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1810 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1811 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1813 Intrinsic::vector_partial_reduce_add)
1816 const auto *FT =
F->getFunctionType();
1818 if (ID == Intrinsic::vector_extract ||
1819 ID == Intrinsic::vector_interleave2)
1822 if (ID != Intrinsic::vector_interleave2)
1824 if (ID == Intrinsic::vector_insert ||
1825 ID == Intrinsic::vector_partial_reduce_add)
1833 if (Name.consume_front(
"reduce.")) {
1835 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1836 if (R.match(Name, &
Groups))
1838 .
Case(
"add", Intrinsic::vector_reduce_add)
1839 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1840 .
Case(
"and", Intrinsic::vector_reduce_and)
1841 .
Case(
"or", Intrinsic::vector_reduce_or)
1842 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1843 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1844 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1845 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1846 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1847 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1848 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1853 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1858 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1859 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1864 auto Args =
F->getFunctionType()->params();
1866 {Args[V2 ? 1 : 0]});
1872 if (Name.consume_front(
"splice"))
1876 if (Name.consume_front(
"experimental.stepvector.")) {
1880 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1885 if (Name.starts_with(
"flt.rounds")) {
1888 Intrinsic::get_rounding);
1893 if (Name.starts_with(
"invariant.group.barrier")) {
1895 auto Args =
F->getFunctionType()->params();
1896 Type* ObjectPtr[1] = {Args[0]};
1899 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1904 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1905 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1906 if (IsLifetimeStart || IsLifetimeEnd) {
1907 if (
F->arg_size() == 2) {
1908 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1909 : Intrinsic::lifetime_end;
1914 F->getArg(1)->getType());
1916 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1936 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1937 .StartsWith(
"memmove.", Intrinsic::memmove)
1939 if (
F->arg_size() == 5) {
1943 F->getFunctionType()->params().slice(0, 3);
1949 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1952 const auto *FT =
F->getFunctionType();
1953 Type *ParamTypes[2] = {
1954 FT->getParamType(0),
1958 Intrinsic::memset, ParamTypes);
1964 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1965 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1966 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1967 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1969 if (MaskedID &&
F->arg_size() == 4) {
1971 if (MaskedID == Intrinsic::masked_load ||
1972 MaskedID == Intrinsic::masked_gather) {
1974 F->getParent(), MaskedID,
1975 {F->getReturnType(), F->getArg(0)->getType()});
1979 F->getParent(), MaskedID,
1980 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1986 if (Name.consume_front(
"nvvm.")) {
1988 if (
F->arg_size() == 1) {
1991 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1992 .Case(
"clz.i", Intrinsic::ctlz)
1993 .
Case(
"popc.i", Intrinsic::ctpop)
1997 {F->getReturnType()});
2000 }
else if (
F->arg_size() == 2) {
2003 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2004 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2005 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2006 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2010 {F->getReturnType()});
2047 F->getParent(), IID,
F->getReturnType(),
2048 F->getFunctionType()->params());
2059 {F->getArg(0)->getType()});
2093 bool Expand =
false;
2094 if (Name.consume_front(
"abs."))
2097 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2098 else if (Name.consume_front(
"fabs."))
2100 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2101 else if (Name.consume_front(
"add."))
2104 else if (Name.consume_front(
"ex2.approx."))
2107 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2108 else if (Name.consume_front(
"atomic.load."))
2117 else if (Name.consume_front(
"atomic."))
2132 else if (Name.consume_front(
"bitcast."))
2135 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2136 else if (Name.consume_front(
"rotate."))
2138 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2139 else if (Name.consume_front(
"ptr.gen.to."))
2142 else if (Name.consume_front(
"ptr."))
2145 else if (Name.consume_front(
"ldg.global."))
2147 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2148 Name.starts_with(
"p."));
2151 .
Case(
"barrier0",
true)
2152 .
Case(
"barrier.n",
true)
2153 .
Case(
"barrier.sync.cnt",
true)
2154 .
Case(
"barrier.sync",
true)
2155 .
Case(
"barrier",
true)
2156 .
Case(
"bar.sync",
true)
2157 .
Case(
"barrier0.popc",
true)
2158 .
Case(
"barrier0.and",
true)
2159 .
Case(
"barrier0.or",
true)
2160 .
Case(
"clz.ll",
true)
2161 .
Case(
"popc.ll",
true)
2163 .
Case(
"swap.lo.hi.b64",
true)
2164 .
Case(
"tanh.approx.f32",
true)
2176 if (Name.starts_with(
"objectsize.")) {
2177 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2178 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2181 Intrinsic::objectsize, Tys);
2188 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2191 F->getParent(), Intrinsic::ptr_annotation,
2192 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2198 if (Name.consume_front(
"riscv.")) {
2201 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2202 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2203 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2204 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2207 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2220 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2221 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2230 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2231 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2232 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2233 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2238 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2247 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2249 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2258 if (Name ==
"stackprotectorcheck") {
2265 if (Name ==
"thread.pointer") {
2267 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2273 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2276 F->getParent(), Intrinsic::var_annotation,
2277 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2280 if (Name.consume_front(
"vector.splice")) {
2281 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2291 if (Name.consume_front(
"wasm.")) {
2294 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2295 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2296 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2301 F->getReturnType());
2305 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2307 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2309 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2328 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2338 std::string
Name =
F->getName().str();
2341 Name,
F->getParent());
2352 if (Result != std::nullopt) {
2368 bool CanUpgradeDebugIntrinsicsToRecords) {
2388 GV->
getName() ==
"llvm.global_dtors")) ||
2403 unsigned N =
Init->getNumOperands();
2404 std::vector<Constant *> NewCtors(
N);
2405 for (
unsigned i = 0; i !=
N; ++i) {
2408 Ctor->getAggregateElement(1),
2422 unsigned NumElts = ResultTy->getNumElements() * 8;
2426 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2436 for (
unsigned l = 0; l != NumElts; l += 16)
2437 for (
unsigned i = 0; i != 16; ++i) {
2438 unsigned Idx = NumElts + i - Shift;
2440 Idx -= NumElts - 16;
2441 Idxs[l + i] = Idx + l;
2444 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2448 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2456 unsigned NumElts = ResultTy->getNumElements() * 8;
2460 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2470 for (
unsigned l = 0; l != NumElts; l += 16)
2471 for (
unsigned i = 0; i != 16; ++i) {
2472 unsigned Idx = i + Shift;
2474 Idx += NumElts - 16;
2475 Idxs[l + i] = Idx + l;
2478 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2482 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2490 Mask = Builder.CreateBitCast(Mask, MaskTy);
2496 for (
unsigned i = 0; i != NumElts; ++i)
2498 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2509 if (
C->isAllOnesValue())
2514 return Builder.CreateSelect(Mask, Op0, Op1);
2521 if (
C->isAllOnesValue())
2525 Mask->getType()->getIntegerBitWidth());
2526 Mask = Builder.CreateBitCast(Mask, MaskTy);
2527 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2528 return Builder.CreateSelect(Mask, Op0, Op1);
2541 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2542 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2547 ShiftVal &= (NumElts - 1);
2556 if (ShiftVal > 16) {
2564 for (
unsigned l = 0; l < NumElts; l += 16) {
2565 for (
unsigned i = 0; i != 16; ++i) {
2566 unsigned Idx = ShiftVal + i;
2567 if (!IsVALIGN && Idx >= 16)
2568 Idx += NumElts - 16;
2569 Indices[l + i] = Idx + l;
2574 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2580 bool ZeroMask,
bool IndexForm) {
2583 unsigned EltWidth = Ty->getScalarSizeInBits();
2584 bool IsFloat = Ty->isFPOrFPVectorTy();
2586 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2587 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2588 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2589 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2590 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2591 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2592 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2593 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2594 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2595 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2596 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2597 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2598 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2599 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2600 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2601 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2602 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2603 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2604 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2605 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2606 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2607 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2608 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2609 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2610 else if (VecWidth == 128 && EltWidth == 16)
2611 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2612 else if (VecWidth == 256 && EltWidth == 16)
2613 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2614 else if (VecWidth == 512 && EltWidth == 16)
2615 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2616 else if (VecWidth == 128 && EltWidth == 8)
2617 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2618 else if (VecWidth == 256 && EltWidth == 8)
2619 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2620 else if (VecWidth == 512 && EltWidth == 8)
2621 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2632 Value *V = Builder.CreateIntrinsic(IID, Args);
2644 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2655 bool IsRotateRight) {
2665 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2666 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2669 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2670 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2715 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2720 bool IsShiftRight,
bool ZeroMask) {
2734 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2735 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2738 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2739 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2754 const Align Alignment =
2756 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2761 if (
C->isAllOnesValue())
2762 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2767 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2773 const Align Alignment =
2782 if (
C->isAllOnesValue())
2783 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2788 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2794 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2795 {Op0, Builder.getInt1(
false)});
2810 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2811 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2812 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2813 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2814 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2817 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2818 LHS = Builder.CreateAnd(
LHS, Mask);
2819 RHS = Builder.CreateAnd(
RHS, Mask);
2836 if (!
C || !
C->isAllOnesValue())
2837 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2842 for (
unsigned i = 0; i != NumElts; ++i)
2844 for (
unsigned i = NumElts; i != 8; ++i)
2845 Indices[i] = NumElts + i % NumElts;
2846 Vec = Builder.CreateShuffleVector(Vec,
2850 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2854 unsigned CC,
bool Signed) {
2862 }
else if (CC == 7) {
2898 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2899 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2901 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2902 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2911 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2917 Name = Name.substr(12);
2922 if (Name.starts_with(
"max.p")) {
2923 if (VecWidth == 128 && EltWidth == 32)
2924 IID = Intrinsic::x86_sse_max_ps;
2925 else if (VecWidth == 128 && EltWidth == 64)
2926 IID = Intrinsic::x86_sse2_max_pd;
2927 else if (VecWidth == 256 && EltWidth == 32)
2928 IID = Intrinsic::x86_avx_max_ps_256;
2929 else if (VecWidth == 256 && EltWidth == 64)
2930 IID = Intrinsic::x86_avx_max_pd_256;
2933 }
else if (Name.starts_with(
"min.p")) {
2934 if (VecWidth == 128 && EltWidth == 32)
2935 IID = Intrinsic::x86_sse_min_ps;
2936 else if (VecWidth == 128 && EltWidth == 64)
2937 IID = Intrinsic::x86_sse2_min_pd;
2938 else if (VecWidth == 256 && EltWidth == 32)
2939 IID = Intrinsic::x86_avx_min_ps_256;
2940 else if (VecWidth == 256 && EltWidth == 64)
2941 IID = Intrinsic::x86_avx_min_pd_256;
2944 }
else if (Name.starts_with(
"pshuf.b.")) {
2945 if (VecWidth == 128)
2946 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2947 else if (VecWidth == 256)
2948 IID = Intrinsic::x86_avx2_pshuf_b;
2949 else if (VecWidth == 512)
2950 IID = Intrinsic::x86_avx512_pshuf_b_512;
2953 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2954 if (VecWidth == 128)
2955 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2956 else if (VecWidth == 256)
2957 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2958 else if (VecWidth == 512)
2959 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2962 }
else if (Name.starts_with(
"pmulh.w.")) {
2963 if (VecWidth == 128)
2964 IID = Intrinsic::x86_sse2_pmulh_w;
2965 else if (VecWidth == 256)
2966 IID = Intrinsic::x86_avx2_pmulh_w;
2967 else if (VecWidth == 512)
2968 IID = Intrinsic::x86_avx512_pmulh_w_512;
2971 }
else if (Name.starts_with(
"pmulhu.w.")) {
2972 if (VecWidth == 128)
2973 IID = Intrinsic::x86_sse2_pmulhu_w;
2974 else if (VecWidth == 256)
2975 IID = Intrinsic::x86_avx2_pmulhu_w;
2976 else if (VecWidth == 512)
2977 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2980 }
else if (Name.starts_with(
"pmaddw.d.")) {
2981 if (VecWidth == 128)
2982 IID = Intrinsic::x86_sse2_pmadd_wd;
2983 else if (VecWidth == 256)
2984 IID = Intrinsic::x86_avx2_pmadd_wd;
2985 else if (VecWidth == 512)
2986 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2989 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2990 if (VecWidth == 128)
2991 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2992 else if (VecWidth == 256)
2993 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2994 else if (VecWidth == 512)
2995 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2998 }
else if (Name.starts_with(
"packsswb.")) {
2999 if (VecWidth == 128)
3000 IID = Intrinsic::x86_sse2_packsswb_128;
3001 else if (VecWidth == 256)
3002 IID = Intrinsic::x86_avx2_packsswb;
3003 else if (VecWidth == 512)
3004 IID = Intrinsic::x86_avx512_packsswb_512;
3007 }
else if (Name.starts_with(
"packssdw.")) {
3008 if (VecWidth == 128)
3009 IID = Intrinsic::x86_sse2_packssdw_128;
3010 else if (VecWidth == 256)
3011 IID = Intrinsic::x86_avx2_packssdw;
3012 else if (VecWidth == 512)
3013 IID = Intrinsic::x86_avx512_packssdw_512;
3016 }
else if (Name.starts_with(
"packuswb.")) {
3017 if (VecWidth == 128)
3018 IID = Intrinsic::x86_sse2_packuswb_128;
3019 else if (VecWidth == 256)
3020 IID = Intrinsic::x86_avx2_packuswb;
3021 else if (VecWidth == 512)
3022 IID = Intrinsic::x86_avx512_packuswb_512;
3025 }
else if (Name.starts_with(
"packusdw.")) {
3026 if (VecWidth == 128)
3027 IID = Intrinsic::x86_sse41_packusdw;
3028 else if (VecWidth == 256)
3029 IID = Intrinsic::x86_avx2_packusdw;
3030 else if (VecWidth == 512)
3031 IID = Intrinsic::x86_avx512_packusdw_512;
3034 }
else if (Name.starts_with(
"vpermilvar.")) {
3035 if (VecWidth == 128 && EltWidth == 32)
3036 IID = Intrinsic::x86_avx_vpermilvar_ps;
3037 else if (VecWidth == 128 && EltWidth == 64)
3038 IID = Intrinsic::x86_avx_vpermilvar_pd;
3039 else if (VecWidth == 256 && EltWidth == 32)
3040 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3041 else if (VecWidth == 256 && EltWidth == 64)
3042 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3043 else if (VecWidth == 512 && EltWidth == 32)
3044 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3045 else if (VecWidth == 512 && EltWidth == 64)
3046 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3049 }
else if (Name ==
"cvtpd2dq.256") {
3050 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3051 }
else if (Name ==
"cvtpd2ps.256") {
3052 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3053 }
else if (Name ==
"cvttpd2dq.256") {
3054 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3055 }
else if (Name ==
"cvttps2dq.128") {
3056 IID = Intrinsic::x86_sse2_cvttps2dq;
3057 }
else if (Name ==
"cvttps2dq.256") {
3058 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3059 }
else if (Name.starts_with(
"permvar.")) {
3061 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3062 IID = Intrinsic::x86_avx2_permps;
3063 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3064 IID = Intrinsic::x86_avx2_permd;
3065 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3066 IID = Intrinsic::x86_avx512_permvar_df_256;
3067 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3068 IID = Intrinsic::x86_avx512_permvar_di_256;
3069 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3070 IID = Intrinsic::x86_avx512_permvar_sf_512;
3071 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3072 IID = Intrinsic::x86_avx512_permvar_si_512;
3073 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3074 IID = Intrinsic::x86_avx512_permvar_df_512;
3075 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3076 IID = Intrinsic::x86_avx512_permvar_di_512;
3077 else if (VecWidth == 128 && EltWidth == 16)
3078 IID = Intrinsic::x86_avx512_permvar_hi_128;
3079 else if (VecWidth == 256 && EltWidth == 16)
3080 IID = Intrinsic::x86_avx512_permvar_hi_256;
3081 else if (VecWidth == 512 && EltWidth == 16)
3082 IID = Intrinsic::x86_avx512_permvar_hi_512;
3083 else if (VecWidth == 128 && EltWidth == 8)
3084 IID = Intrinsic::x86_avx512_permvar_qi_128;
3085 else if (VecWidth == 256 && EltWidth == 8)
3086 IID = Intrinsic::x86_avx512_permvar_qi_256;
3087 else if (VecWidth == 512 && EltWidth == 8)
3088 IID = Intrinsic::x86_avx512_permvar_qi_512;
3091 }
else if (Name.starts_with(
"dbpsadbw.")) {
3092 if (VecWidth == 128)
3093 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3094 else if (VecWidth == 256)
3095 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3096 else if (VecWidth == 512)
3097 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3100 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3101 if (VecWidth == 128)
3102 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3103 else if (VecWidth == 256)
3104 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3105 else if (VecWidth == 512)
3106 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3109 }
else if (Name.starts_with(
"conflict.")) {
3110 if (Name[9] ==
'd' && VecWidth == 128)
3111 IID = Intrinsic::x86_avx512_conflict_d_128;
3112 else if (Name[9] ==
'd' && VecWidth == 256)
3113 IID = Intrinsic::x86_avx512_conflict_d_256;
3114 else if (Name[9] ==
'd' && VecWidth == 512)
3115 IID = Intrinsic::x86_avx512_conflict_d_512;
3116 else if (Name[9] ==
'q' && VecWidth == 128)
3117 IID = Intrinsic::x86_avx512_conflict_q_128;
3118 else if (Name[9] ==
'q' && VecWidth == 256)
3119 IID = Intrinsic::x86_avx512_conflict_q_256;
3120 else if (Name[9] ==
'q' && VecWidth == 512)
3121 IID = Intrinsic::x86_avx512_conflict_q_512;
3124 }
else if (Name.starts_with(
"pavg.")) {
3125 if (Name[5] ==
'b' && VecWidth == 128)
3126 IID = Intrinsic::x86_sse2_pavg_b;
3127 else if (Name[5] ==
'b' && VecWidth == 256)
3128 IID = Intrinsic::x86_avx2_pavg_b;
3129 else if (Name[5] ==
'b' && VecWidth == 512)
3130 IID = Intrinsic::x86_avx512_pavg_b_512;
3131 else if (Name[5] ==
'w' && VecWidth == 128)
3132 IID = Intrinsic::x86_sse2_pavg_w;
3133 else if (Name[5] ==
'w' && VecWidth == 256)
3134 IID = Intrinsic::x86_avx2_pavg_w;
3135 else if (Name[5] ==
'w' && VecWidth == 512)
3136 IID = Intrinsic::x86_avx512_pavg_w_512;
3145 Rep = Builder.CreateIntrinsic(IID, Args);
3156 if (AsmStr->find(
"mov\tfp") == 0 &&
3157 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3158 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3159 AsmStr->replace(Pos, 1,
";");
3165 Value *Rep =
nullptr;
3167 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3169 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3170 {Arg, Builder.getTrue()},
3172 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3173 Type *Ty = (Name ==
"abs.bf16")
3177 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3178 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3179 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3180 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3181 : Intrinsic::nvvm_fabs;
3182 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3183 }
else if (Name.consume_front(
"add.")) {
3186 assert(
FAdd &&
"unsupported nvvm.add.* intrinsic");
3189 Rep = Builder.CreateIntrinsic(
3191 {A, CI->getArgOperand(1),
3192 Builder.getInt32(static_cast<int>(RoundingMode))});
3193 }
else if (Name.consume_front(
"ex2.approx.")) {
3195 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3196 : Intrinsic::nvvm_ex2_approx;
3197 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3198 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3199 Name.starts_with(
"atomic.load.add.f64.p")) {
3202 Rep = Builder.CreateAtomicRMW(
3208 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3209 Name.starts_with(
"atomic.load.dec.32.p")) {
3214 Rep = Builder.CreateAtomicRMW(
3218 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3224 Op.contains(
".cta.") ?
"block" :
"");
3225 if (
Op.starts_with(
"cas.")) {
3227 Value *Pair = Builder.CreateAtomicCmpXchg(
3230 Rep = Builder.CreateExtractValue(Pair, 0);
3248 "unexpected nvvm scoped atomic intrinsic");
3249 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3252 }
else if (Name ==
"clz.ll") {
3255 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3256 {Arg, Builder.getFalse()},
3258 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3259 }
else if (Name ==
"popc.ll") {
3263 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3264 Arg,
nullptr,
"ctpop");
3265 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3266 }
else if (Name ==
"h2f") {
3268 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3269 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3270 }
else if (Name.consume_front(
"bitcast.") &&
3271 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3274 }
else if (Name ==
"rotate.b32") {
3277 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3278 {Arg, Arg, ShiftAmt});
3279 }
else if (Name ==
"rotate.b64") {
3283 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3284 {Arg, Arg, ZExtShiftAmt});
3285 }
else if (Name ==
"rotate.right.b64") {
3289 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3290 {Arg, Arg, ZExtShiftAmt});
3291 }
else if (Name ==
"swap.lo.hi.b64") {
3294 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3295 {Arg, Arg, Builder.getInt64(32)});
3296 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3299 Name.starts_with(
".to.gen"))) {
3301 }
else if (Name.consume_front(
"ldg.global")) {
3305 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3308 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3310 }
else if (Name ==
"tanh.approx.f32") {
3314 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3316 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3318 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3319 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3321 }
else if (Name ==
"barrier") {
3322 Rep = Builder.CreateIntrinsic(
3323 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3325 }
else if (Name ==
"barrier.sync") {
3326 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3328 }
else if (Name ==
"barrier.sync.cnt") {
3329 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3331 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3332 Name ==
"barrier0.or") {
3334 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3338 .
Case(
"barrier0.popc",
3339 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3340 .
Case(
"barrier0.and",
3341 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3342 .
Case(
"barrier0.or",
3343 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3344 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3345 Rep = Builder.CreateZExt(Bar, CI->
getType());
3359 ? Builder.CreateBitCast(Arg, NewType)
3362 Rep = Builder.CreateCall(NewFn, Args);
3363 if (
F->getReturnType()->isIntegerTy())
3364 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3374 Value *Rep =
nullptr;
3376 if (Name.starts_with(
"sse4a.movnt.")) {
3388 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3391 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3392 }
else if (Name.starts_with(
"avx.movnt.") ||
3393 Name.starts_with(
"avx512.storent.")) {
3405 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3406 }
else if (Name ==
"sse2.storel.dq") {
3411 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3412 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3413 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3414 }
else if (Name.starts_with(
"sse.storeu.") ||
3415 Name.starts_with(
"sse2.storeu.") ||
3416 Name.starts_with(
"avx.storeu.")) {
3419 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3420 }
else if (Name ==
"avx512.mask.store.ss") {
3424 }
else if (Name.starts_with(
"avx512.mask.store")) {
3426 bool Aligned = Name[17] !=
'u';
3429 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3432 bool CmpEq = Name[9] ==
'e';
3435 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3436 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3443 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3444 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3446 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3447 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3448 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3449 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3450 Name.starts_with(
"sse2.sqrt.p") ||
3451 Name.starts_with(
"sse.sqrt.p")) {
3452 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3453 {CI->getArgOperand(0)});
3454 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3458 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3459 : Intrinsic::x86_avx512_sqrt_pd_512;
3462 Rep = Builder.CreateIntrinsic(IID, Args);
3464 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3465 {CI->getArgOperand(0)});
3469 }
else if (Name.starts_with(
"avx512.ptestm") ||
3470 Name.starts_with(
"avx512.ptestnm")) {
3474 Rep = Builder.CreateAnd(Op0, Op1);
3480 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3482 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3485 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3488 }
else if (Name.starts_with(
"avx512.kunpck")) {
3493 for (
unsigned i = 0; i != NumElts; ++i)
3502 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3503 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3504 }
else if (Name ==
"avx512.kand.w") {
3507 Rep = Builder.CreateAnd(
LHS,
RHS);
3508 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3509 }
else if (Name ==
"avx512.kandn.w") {
3512 LHS = Builder.CreateNot(
LHS);
3513 Rep = Builder.CreateAnd(
LHS,
RHS);
3514 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3515 }
else if (Name ==
"avx512.kor.w") {
3518 Rep = Builder.CreateOr(
LHS,
RHS);
3519 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3520 }
else if (Name ==
"avx512.kxor.w") {
3523 Rep = Builder.CreateXor(
LHS,
RHS);
3524 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3525 }
else if (Name ==
"avx512.kxnor.w") {
3528 LHS = Builder.CreateNot(
LHS);
3529 Rep = Builder.CreateXor(
LHS,
RHS);
3530 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3531 }
else if (Name ==
"avx512.knot.w") {
3533 Rep = Builder.CreateNot(Rep);
3534 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3535 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3538 Rep = Builder.CreateOr(
LHS,
RHS);
3539 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3541 if (Name[14] ==
'c')
3545 Rep = Builder.CreateICmpEQ(Rep,
C);
3546 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3547 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3548 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3549 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3550 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3553 ConstantInt::get(I32Ty, 0));
3555 ConstantInt::get(I32Ty, 0));
3557 if (Name.contains(
".add."))
3558 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3559 else if (Name.contains(
".sub."))
3560 EltOp = Builder.CreateFSub(Elt0, Elt1);
3561 else if (Name.contains(
".mul."))
3562 EltOp = Builder.CreateFMul(Elt0, Elt1);
3564 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3565 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3566 ConstantInt::get(I32Ty, 0));
3567 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3569 bool CmpEq = Name[16] ==
'e';
3571 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3573 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3580 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3583 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3586 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3593 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3595 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3596 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3598 if (VecWidth == 128 && EltWidth == 32)
3599 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3600 else if (VecWidth == 256 && EltWidth == 32)
3601 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3602 else if (VecWidth == 512 && EltWidth == 32)
3603 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3604 else if (VecWidth == 128 && EltWidth == 64)
3605 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3606 else if (VecWidth == 256 && EltWidth == 64)
3607 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3608 else if (VecWidth == 512 && EltWidth == 64)
3609 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3616 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3619 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3620 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3622 if (VecWidth == 128 && EltWidth == 32)
3623 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3624 else if (VecWidth == 256 && EltWidth == 32)
3625 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3626 else if (VecWidth == 512 && EltWidth == 32)
3627 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3628 else if (VecWidth == 128 && EltWidth == 64)
3629 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3630 else if (VecWidth == 256 && EltWidth == 64)
3631 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3632 else if (VecWidth == 512 && EltWidth == 64)
3633 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3638 if (VecWidth == 512)
3640 Args.push_back(Mask);
3642 Rep = Builder.CreateIntrinsic(IID, Args);
3643 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3647 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3650 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3651 Name.starts_with(
"avx512.cvtw2mask.") ||
3652 Name.starts_with(
"avx512.cvtd2mask.") ||
3653 Name.starts_with(
"avx512.cvtq2mask.")) {
3658 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3659 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3660 Name.starts_with(
"avx512.mask.pabs")) {
3662 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3663 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3664 Name.starts_with(
"avx512.mask.pmaxs")) {
3666 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3667 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3668 Name.starts_with(
"avx512.mask.pmaxu")) {
3670 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3671 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3672 Name.starts_with(
"avx512.mask.pmins")) {
3674 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3675 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3676 Name.starts_with(
"avx512.mask.pminu")) {
3678 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3679 Name ==
"avx512.pmulu.dq.512" ||
3680 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3682 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3683 Name ==
"avx512.pmul.dq.512" ||
3684 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3686 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3687 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3692 }
else if (Name ==
"avx512.cvtusi2sd") {
3697 }
else if (Name ==
"sse2.cvtss2sd") {
3699 Rep = Builder.CreateFPExt(
3702 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3703 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3704 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3705 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3706 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3707 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3708 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3709 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3710 Name ==
"avx512.mask.cvtqq2ps.256" ||
3711 Name ==
"avx512.mask.cvtqq2ps.512" ||
3712 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3713 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3714 Name ==
"avx.cvt.ps2.pd.256" ||
3715 Name ==
"avx512.mask.cvtps2pd.128" ||
3716 Name ==
"avx512.mask.cvtps2pd.256") {
3721 unsigned NumDstElts = DstTy->getNumElements();
3722 if (NumDstElts < SrcTy->getNumElements()) {
3723 assert(NumDstElts == 2 &&
"Unexpected vector size");
3724 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3727 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3728 bool IsUnsigned = Name.contains(
"cvtu");
3730 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3734 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3735 : Intrinsic::x86_avx512_sitofp_round;
3736 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3739 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3740 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3746 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3747 Name.starts_with(
"vcvtph2ps.")) {
3751 unsigned NumDstElts = DstTy->getNumElements();
3752 if (NumDstElts != SrcTy->getNumElements()) {
3753 assert(NumDstElts == 4 &&
"Unexpected vector size");
3754 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3756 Rep = Builder.CreateBitCast(
3758 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3762 }
else if (Name.starts_with(
"avx512.mask.load")) {
3764 bool Aligned = Name[16] !=
'u';
3767 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3771 ResultTy->getNumElements());
3772 Rep = Builder.CreateIntrinsic(
3773 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3775 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3781 Rep = Builder.CreateIntrinsic(
3782 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3784 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3785 Name.starts_with(
"avx512.mask.expand.")) {
3789 ResultTy->getNumElements());
3791 bool IsCompress = Name[12] ==
'c';
3792 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3793 : Intrinsic::x86_avx512_mask_expand;
3794 Rep = Builder.CreateIntrinsic(
3796 }
else if (Name.starts_with(
"xop.vpcom")) {
3798 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3799 Name.ends_with(
"uq"))
3801 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3802 Name.ends_with(
"d") || Name.ends_with(
"q"))
3811 Name = Name.substr(9);
3812 if (Name.starts_with(
"lt"))
3814 else if (Name.starts_with(
"le"))
3816 else if (Name.starts_with(
"gt"))
3818 else if (Name.starts_with(
"ge"))
3820 else if (Name.starts_with(
"eq"))
3822 else if (Name.starts_with(
"ne"))
3824 else if (Name.starts_with(
"false"))
3826 else if (Name.starts_with(
"true"))
3833 }
else if (Name.starts_with(
"xop.vpcmov")) {
3835 Value *NotSel = Builder.CreateNot(Sel);
3838 Rep = Builder.CreateOr(Sel0, Sel1);
3839 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3840 Name.starts_with(
"avx512.mask.prol")) {
3842 }
else if (Name.starts_with(
"avx512.pror") ||
3843 Name.starts_with(
"avx512.mask.pror")) {
3845 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3846 Name.starts_with(
"avx512.mask.vpshld") ||
3847 Name.starts_with(
"avx512.maskz.vpshld")) {
3848 bool ZeroMask = Name[11] ==
'z';
3850 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3851 Name.starts_with(
"avx512.mask.vpshrd") ||
3852 Name.starts_with(
"avx512.maskz.vpshrd")) {
3853 bool ZeroMask = Name[11] ==
'z';
3855 }
else if (Name ==
"sse42.crc32.64.8") {
3858 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3860 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3861 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3862 Name.starts_with(
"avx512.vbroadcast.s")) {
3865 Type *EltTy = VecTy->getElementType();
3866 unsigned EltNum = VecTy->getNumElements();
3870 for (
unsigned I = 0;
I < EltNum; ++
I)
3871 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3872 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3873 Name.starts_with(
"sse41.pmovzx") ||
3874 Name.starts_with(
"avx2.pmovsx") ||
3875 Name.starts_with(
"avx2.pmovzx") ||
3876 Name.starts_with(
"avx512.mask.pmovsx") ||
3877 Name.starts_with(
"avx512.mask.pmovzx")) {
3879 unsigned NumDstElts = DstTy->getNumElements();
3883 for (
unsigned i = 0; i != NumDstElts; ++i)
3888 bool DoSext = Name.contains(
"pmovsx");
3890 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3895 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3896 Name ==
"avx512.mask.pmov.qd.512" ||
3897 Name ==
"avx512.mask.pmov.wb.256" ||
3898 Name ==
"avx512.mask.pmov.wb.512") {
3903 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3904 Name ==
"avx2.vbroadcasti128") {
3910 if (NumSrcElts == 2)
3913 Rep = Builder.CreateShuffleVector(
Load,
3915 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3916 Name.starts_with(
"avx512.mask.shuf.f")) {
3921 unsigned ControlBitsMask = NumLanes - 1;
3922 unsigned NumControlBits = NumLanes / 2;
3925 for (
unsigned l = 0; l != NumLanes; ++l) {
3926 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3928 if (l >= NumLanes / 2)
3929 LaneMask += NumLanes;
3930 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3931 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3937 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3938 Name.starts_with(
"avx512.mask.broadcasti")) {
3941 unsigned NumDstElts =
3945 for (
unsigned i = 0; i != NumDstElts; ++i)
3946 ShuffleMask[i] = i % NumSrcElts;
3952 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3953 Name.starts_with(
"avx2.vbroadcast") ||
3954 Name.starts_with(
"avx512.pbroadcast") ||
3955 Name.starts_with(
"avx512.mask.broadcast.s")) {
3962 Rep = Builder.CreateShuffleVector(
Op, M);
3967 }
else if (Name.starts_with(
"sse2.padds.") ||
3968 Name.starts_with(
"avx2.padds.") ||
3969 Name.starts_with(
"avx512.padds.") ||
3970 Name.starts_with(
"avx512.mask.padds.")) {
3972 }
else if (Name.starts_with(
"sse2.psubs.") ||
3973 Name.starts_with(
"avx2.psubs.") ||
3974 Name.starts_with(
"avx512.psubs.") ||
3975 Name.starts_with(
"avx512.mask.psubs.")) {
3977 }
else if (Name.starts_with(
"sse2.paddus.") ||
3978 Name.starts_with(
"avx2.paddus.") ||
3979 Name.starts_with(
"avx512.mask.paddus.")) {
3981 }
else if (Name.starts_with(
"sse2.psubus.") ||
3982 Name.starts_with(
"avx2.psubus.") ||
3983 Name.starts_with(
"avx512.mask.psubus.")) {
3985 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3990 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3994 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3999 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4004 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4005 Name ==
"avx512.psll.dq.512") {
4009 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4010 Name ==
"avx512.psrl.dq.512") {
4014 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4015 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4016 Name.starts_with(
"avx2.pblendd.")) {
4021 unsigned NumElts = VecTy->getNumElements();
4024 for (
unsigned i = 0; i != NumElts; ++i)
4025 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4027 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4028 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4029 Name ==
"avx2.vinserti128" ||
4030 Name.starts_with(
"avx512.mask.insert")) {
4034 unsigned DstNumElts =
4036 unsigned SrcNumElts =
4038 unsigned Scale = DstNumElts / SrcNumElts;
4045 for (
unsigned i = 0; i != SrcNumElts; ++i)
4047 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4048 Idxs[i] = SrcNumElts;
4049 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4063 for (
unsigned i = 0; i != DstNumElts; ++i)
4066 for (
unsigned i = 0; i != SrcNumElts; ++i)
4067 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4068 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4074 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4075 Name ==
"avx2.vextracti128" ||
4076 Name.starts_with(
"avx512.mask.vextract")) {
4079 unsigned DstNumElts =
4081 unsigned SrcNumElts =
4083 unsigned Scale = SrcNumElts / DstNumElts;
4090 for (
unsigned i = 0; i != DstNumElts; ++i) {
4091 Idxs[i] = i + (
Imm * DstNumElts);
4093 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4099 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4100 Name.starts_with(
"avx512.mask.perm.di.")) {
4104 unsigned NumElts = VecTy->getNumElements();
4107 for (
unsigned i = 0; i != NumElts; ++i)
4108 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4110 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4115 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4127 unsigned HalfSize = NumElts / 2;
4139 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4140 for (
unsigned i = 0; i < HalfSize; ++i)
4141 ShuffleMask[i] = StartIndex + i;
4144 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4145 for (
unsigned i = 0; i < HalfSize; ++i)
4146 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4148 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4150 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4151 Name.starts_with(
"avx512.mask.vpermil.p") ||
4152 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4156 unsigned NumElts = VecTy->getNumElements();
4158 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4159 unsigned IdxMask = ((1 << IdxSize) - 1);
4165 for (
unsigned i = 0; i != NumElts; ++i)
4166 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4168 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4173 }
else if (Name ==
"sse2.pshufl.w" ||
4174 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4179 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4183 for (
unsigned l = 0; l != NumElts; l += 8) {
4184 for (
unsigned i = 0; i != 4; ++i)
4185 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4186 for (
unsigned i = 4; i != 8; ++i)
4187 Idxs[i + l] = i + l;
4190 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4195 }
else if (Name ==
"sse2.pshufh.w" ||
4196 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4201 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4205 for (
unsigned l = 0; l != NumElts; l += 8) {
4206 for (
unsigned i = 0; i != 4; ++i)
4207 Idxs[i + l] = i + l;
4208 for (
unsigned i = 0; i != 4; ++i)
4209 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4212 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4217 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4224 unsigned HalfLaneElts = NumLaneElts / 2;
4227 for (
unsigned i = 0; i != NumElts; ++i) {
4229 Idxs[i] = i - (i % NumLaneElts);
4231 if ((i % NumLaneElts) >= HalfLaneElts)
4235 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4238 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4242 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4243 Name.starts_with(
"avx512.mask.movshdup") ||
4244 Name.starts_with(
"avx512.mask.movsldup")) {
4250 if (Name.starts_with(
"avx512.mask.movshdup."))
4254 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4255 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4256 Idxs[i + l + 0] = i + l +
Offset;
4257 Idxs[i + l + 1] = i + l +
Offset;
4260 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4264 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4265 Name.starts_with(
"avx512.mask.unpckl.")) {
4272 for (
int l = 0; l != NumElts; l += NumLaneElts)
4273 for (
int i = 0; i != NumLaneElts; ++i)
4274 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4276 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4280 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4281 Name.starts_with(
"avx512.mask.unpckh.")) {
4288 for (
int l = 0; l != NumElts; l += NumLaneElts)
4289 for (
int i = 0; i != NumLaneElts; ++i)
4290 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4292 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4296 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4297 Name.starts_with(
"avx512.mask.pand.")) {
4300 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4302 Rep = Builder.CreateBitCast(Rep, FTy);
4305 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4306 Name.starts_with(
"avx512.mask.pandn.")) {
4309 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4310 Rep = Builder.CreateAnd(Rep,
4312 Rep = Builder.CreateBitCast(Rep, FTy);
4315 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4316 Name.starts_with(
"avx512.mask.por.")) {
4319 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4321 Rep = Builder.CreateBitCast(Rep, FTy);
4324 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4325 Name.starts_with(
"avx512.mask.pxor.")) {
4328 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4330 Rep = Builder.CreateBitCast(Rep, FTy);
4333 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4337 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4341 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4345 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4346 if (Name.ends_with(
".512")) {
4348 if (Name[17] ==
's')
4349 IID = Intrinsic::x86_avx512_add_ps_512;
4351 IID = Intrinsic::x86_avx512_add_pd_512;
4353 Rep = Builder.CreateIntrinsic(
4361 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4362 if (Name.ends_with(
".512")) {
4364 if (Name[17] ==
's')
4365 IID = Intrinsic::x86_avx512_div_ps_512;
4367 IID = Intrinsic::x86_avx512_div_pd_512;
4369 Rep = Builder.CreateIntrinsic(
4377 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4378 if (Name.ends_with(
".512")) {
4380 if (Name[17] ==
's')
4381 IID = Intrinsic::x86_avx512_mul_ps_512;
4383 IID = Intrinsic::x86_avx512_mul_pd_512;
4385 Rep = Builder.CreateIntrinsic(
4393 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4394 if (Name.ends_with(
".512")) {
4396 if (Name[17] ==
's')
4397 IID = Intrinsic::x86_avx512_sub_ps_512;
4399 IID = Intrinsic::x86_avx512_sub_pd_512;
4401 Rep = Builder.CreateIntrinsic(
4409 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4410 Name.starts_with(
"avx512.mask.min.p")) &&
4411 Name.drop_front(18) ==
".512") {
4412 bool IsDouble = Name[17] ==
'd';
4413 bool IsMin = Name[13] ==
'i';
4415 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4416 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4419 Rep = Builder.CreateIntrinsic(
4424 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4426 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4427 {CI->getArgOperand(0), Builder.getInt1(false)});
4430 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4431 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4432 bool IsVariable = Name[16] ==
'v';
4433 char Size = Name[16] ==
'.' ? Name[17]
4434 : Name[17] ==
'.' ? Name[18]
4435 : Name[18] ==
'.' ? Name[19]
4439 if (IsVariable && Name[17] !=
'.') {
4440 if (
Size ==
'd' && Name[17] ==
'2')
4441 IID = Intrinsic::x86_avx2_psllv_q;
4442 else if (
Size ==
'd' && Name[17] ==
'4')
4443 IID = Intrinsic::x86_avx2_psllv_q_256;
4444 else if (
Size ==
's' && Name[17] ==
'4')
4445 IID = Intrinsic::x86_avx2_psllv_d;
4446 else if (
Size ==
's' && Name[17] ==
'8')
4447 IID = Intrinsic::x86_avx2_psllv_d_256;
4448 else if (
Size ==
'h' && Name[17] ==
'8')
4449 IID = Intrinsic::x86_avx512_psllv_w_128;
4450 else if (
Size ==
'h' && Name[17] ==
'1')
4451 IID = Intrinsic::x86_avx512_psllv_w_256;
4452 else if (Name[17] ==
'3' && Name[18] ==
'2')
4453 IID = Intrinsic::x86_avx512_psllv_w_512;
4456 }
else if (Name.ends_with(
".128")) {
4458 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4459 : Intrinsic::x86_sse2_psll_d;
4460 else if (
Size ==
'q')
4461 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4462 : Intrinsic::x86_sse2_psll_q;
4463 else if (
Size ==
'w')
4464 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4465 : Intrinsic::x86_sse2_psll_w;
4468 }
else if (Name.ends_with(
".256")) {
4470 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4471 : Intrinsic::x86_avx2_psll_d;
4472 else if (
Size ==
'q')
4473 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4474 : Intrinsic::x86_avx2_psll_q;
4475 else if (
Size ==
'w')
4476 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4477 : Intrinsic::x86_avx2_psll_w;
4482 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4483 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4484 : Intrinsic::x86_avx512_psll_d_512;
4485 else if (
Size ==
'q')
4486 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4487 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4488 : Intrinsic::x86_avx512_psll_q_512;
4489 else if (
Size ==
'w')
4490 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4491 : Intrinsic::x86_avx512_psll_w_512;
4497 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
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_psrlv_q;
4509 else if (
Size ==
'd' && Name[17] ==
'4')
4510 IID = Intrinsic::x86_avx2_psrlv_q_256;
4511 else if (
Size ==
's' && Name[17] ==
'4')
4512 IID = Intrinsic::x86_avx2_psrlv_d;
4513 else if (
Size ==
's' && Name[17] ==
'8')
4514 IID = Intrinsic::x86_avx2_psrlv_d_256;
4515 else if (
Size ==
'h' && Name[17] ==
'8')
4516 IID = Intrinsic::x86_avx512_psrlv_w_128;
4517 else if (
Size ==
'h' && Name[17] ==
'1')
4518 IID = Intrinsic::x86_avx512_psrlv_w_256;
4519 else if (Name[17] ==
'3' && Name[18] ==
'2')
4520 IID = Intrinsic::x86_avx512_psrlv_w_512;
4523 }
else if (Name.ends_with(
".128")) {
4525 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4526 : Intrinsic::x86_sse2_psrl_d;
4527 else if (
Size ==
'q')
4528 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4529 : Intrinsic::x86_sse2_psrl_q;
4530 else if (
Size ==
'w')
4531 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4532 : Intrinsic::x86_sse2_psrl_w;
4535 }
else if (Name.ends_with(
".256")) {
4537 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4538 : Intrinsic::x86_avx2_psrl_d;
4539 else if (
Size ==
'q')
4540 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4541 : Intrinsic::x86_avx2_psrl_q;
4542 else if (
Size ==
'w')
4543 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4544 : Intrinsic::x86_avx2_psrl_w;
4549 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4550 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4551 : Intrinsic::x86_avx512_psrl_d_512;
4552 else if (
Size ==
'q')
4553 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4554 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4555 : Intrinsic::x86_avx512_psrl_q_512;
4556 else if (
Size ==
'w')
4557 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4558 : Intrinsic::x86_avx512_psrl_w_512;
4564 }
else if (Name.starts_with(
"avx512.mask.psra")) {
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 ==
's' && Name[17] ==
'4')
4575 IID = Intrinsic::x86_avx2_psrav_d;
4576 else if (
Size ==
's' && Name[17] ==
'8')
4577 IID = Intrinsic::x86_avx2_psrav_d_256;
4578 else if (
Size ==
'h' && Name[17] ==
'8')
4579 IID = Intrinsic::x86_avx512_psrav_w_128;
4580 else if (
Size ==
'h' && Name[17] ==
'1')
4581 IID = Intrinsic::x86_avx512_psrav_w_256;
4582 else if (Name[17] ==
'3' && Name[18] ==
'2')
4583 IID = Intrinsic::x86_avx512_psrav_w_512;
4586 }
else if (Name.ends_with(
".128")) {
4588 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4589 : Intrinsic::x86_sse2_psra_d;
4590 else if (
Size ==
'q')
4591 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4592 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4593 : Intrinsic::x86_avx512_psra_q_128;
4594 else if (
Size ==
'w')
4595 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4596 : Intrinsic::x86_sse2_psra_w;
4599 }
else if (Name.ends_with(
".256")) {
4601 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4602 : Intrinsic::x86_avx2_psra_d;
4603 else if (
Size ==
'q')
4604 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4605 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4606 : Intrinsic::x86_avx512_psra_q_256;
4607 else if (
Size ==
'w')
4608 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4609 : Intrinsic::x86_avx2_psra_w;
4614 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4615 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4616 : Intrinsic::x86_avx512_psra_d_512;
4617 else if (
Size ==
'q')
4618 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4619 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4620 : Intrinsic::x86_avx512_psra_q_512;
4621 else if (
Size ==
'w')
4622 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4623 : Intrinsic::x86_avx512_psra_w_512;
4629 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4631 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4633 }
else if (Name.ends_with(
".movntdqa")) {
4637 LoadInst *LI = Builder.CreateAlignedLoad(
4642 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4643 Name.starts_with(
"fma.vfmsub.") ||
4644 Name.starts_with(
"fma.vfnmadd.") ||
4645 Name.starts_with(
"fma.vfnmsub.")) {
4646 bool NegMul = Name[6] ==
'n';
4647 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4648 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4659 if (NegMul && !IsScalar)
4660 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4661 if (NegMul && IsScalar)
4662 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4664 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4666 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4670 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4678 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4682 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4683 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4684 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4685 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4686 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4687 bool IsMask3 = Name[11] ==
'3';
4688 bool IsMaskZ = Name[11] ==
'z';
4690 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4691 bool NegMul = Name[2] ==
'n';
4692 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4698 if (NegMul && (IsMask3 || IsMaskZ))
4699 A = Builder.CreateFNeg(
A);
4700 if (NegMul && !(IsMask3 || IsMaskZ))
4701 B = Builder.CreateFNeg(
B);
4703 C = Builder.CreateFNeg(
C);
4705 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4706 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4707 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4714 if (Name.back() ==
'd')
4715 IID = Intrinsic::x86_avx512_vfmadd_f64;
4717 IID = Intrinsic::x86_avx512_vfmadd_f32;
4718 Rep = Builder.CreateIntrinsic(IID,
Ops);
4720 Rep = Builder.CreateFMA(
A,
B,
C);
4729 if (NegAcc && IsMask3)
4734 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4736 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4737 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4738 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4739 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4740 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4741 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4742 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4743 bool IsMask3 = Name[11] ==
'3';
4744 bool IsMaskZ = Name[11] ==
'z';
4746 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4747 bool NegMul = Name[2] ==
'n';
4748 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4754 if (NegMul && (IsMask3 || IsMaskZ))
4755 A = Builder.CreateFNeg(
A);
4756 if (NegMul && !(IsMask3 || IsMaskZ))
4757 B = Builder.CreateFNeg(
B);
4759 C = Builder.CreateFNeg(
C);
4766 if (Name[Name.size() - 5] ==
's')
4767 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4769 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4773 Rep = Builder.CreateFMA(
A,
B,
C);
4781 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4785 if (VecWidth == 128 && EltWidth == 32)
4786 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4787 else if (VecWidth == 256 && EltWidth == 32)
4788 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4789 else if (VecWidth == 128 && EltWidth == 64)
4790 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4791 else if (VecWidth == 256 && EltWidth == 64)
4792 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4798 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4799 Rep = Builder.CreateIntrinsic(IID,
Ops);
4800 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4801 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4802 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4803 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4804 bool IsMask3 = Name[11] ==
'3';
4805 bool IsMaskZ = Name[11] ==
'z';
4807 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4808 bool IsSubAdd = Name[3] ==
's';
4812 if (Name[Name.size() - 5] ==
's')
4813 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4815 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4820 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4822 Rep = Builder.CreateIntrinsic(IID,
Ops);
4831 Value *Odd = Builder.CreateCall(FMA,
Ops);
4832 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4833 Value *Even = Builder.CreateCall(FMA,
Ops);
4839 for (
int i = 0; i != NumElts; ++i)
4840 Idxs[i] = i + (i % 2) * NumElts;
4842 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4850 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4851 Name.starts_with(
"avx512.maskz.pternlog.")) {
4852 bool ZeroMask = Name[11] ==
'z';
4856 if (VecWidth == 128 && EltWidth == 32)
4857 IID = Intrinsic::x86_avx512_pternlog_d_128;
4858 else if (VecWidth == 256 && EltWidth == 32)
4859 IID = Intrinsic::x86_avx512_pternlog_d_256;
4860 else if (VecWidth == 512 && EltWidth == 32)
4861 IID = Intrinsic::x86_avx512_pternlog_d_512;
4862 else if (VecWidth == 128 && EltWidth == 64)
4863 IID = Intrinsic::x86_avx512_pternlog_q_128;
4864 else if (VecWidth == 256 && EltWidth == 64)
4865 IID = Intrinsic::x86_avx512_pternlog_q_256;
4866 else if (VecWidth == 512 && EltWidth == 64)
4867 IID = Intrinsic::x86_avx512_pternlog_q_512;
4873 Rep = Builder.CreateIntrinsic(IID, Args);
4877 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4878 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4879 bool ZeroMask = Name[11] ==
'z';
4880 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4883 if (VecWidth == 128 && !
High)
4884 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4885 else if (VecWidth == 256 && !
High)
4886 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4887 else if (VecWidth == 512 && !
High)
4888 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4889 else if (VecWidth == 128 &&
High)
4890 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4891 else if (VecWidth == 256 &&
High)
4892 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4893 else if (VecWidth == 512 &&
High)
4894 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4900 Rep = Builder.CreateIntrinsic(IID, Args);
4904 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4905 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4906 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4907 bool ZeroMask = Name[11] ==
'z';
4908 bool IndexForm = Name[17] ==
'i';
4910 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4911 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4912 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4913 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4914 bool ZeroMask = Name[11] ==
'z';
4915 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4918 if (VecWidth == 128 && !IsSaturating)
4919 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4920 else if (VecWidth == 256 && !IsSaturating)
4921 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4922 else if (VecWidth == 512 && !IsSaturating)
4923 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4924 else if (VecWidth == 128 && IsSaturating)
4925 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4926 else if (VecWidth == 256 && IsSaturating)
4927 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4928 else if (VecWidth == 512 && IsSaturating)
4929 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4939 if (Args[1]->
getType()->isVectorTy() &&
4942 ->isIntegerTy(32) &&
4943 Args[2]->
getType()->isVectorTy() &&
4946 ->isIntegerTy(32)) {
4947 Type *NewArgType =
nullptr;
4948 if (VecWidth == 128)
4950 else if (VecWidth == 256)
4952 else if (VecWidth == 512)
4958 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4959 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4962 Rep = Builder.CreateIntrinsic(IID, Args);
4966 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4967 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4968 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4969 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4970 bool ZeroMask = Name[11] ==
'z';
4971 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4974 if (VecWidth == 128 && !IsSaturating)
4975 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4976 else if (VecWidth == 256 && !IsSaturating)
4977 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4978 else if (VecWidth == 512 && !IsSaturating)
4979 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4980 else if (VecWidth == 128 && IsSaturating)
4981 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4982 else if (VecWidth == 256 && IsSaturating)
4983 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4984 else if (VecWidth == 512 && IsSaturating)
4985 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4995 if (Args[1]->
getType()->isVectorTy() &&
4998 ->isIntegerTy(32) &&
4999 Args[2]->
getType()->isVectorTy() &&
5002 ->isIntegerTy(32)) {
5003 Type *NewArgType =
nullptr;
5004 if (VecWidth == 128)
5006 else if (VecWidth == 256)
5008 else if (VecWidth == 512)
5014 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5015 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5018 Rep = Builder.CreateIntrinsic(IID, Args);
5022 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5023 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5024 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5026 if (Name[0] ==
'a' && Name.back() ==
'2')
5027 IID = Intrinsic::x86_addcarry_32;
5028 else if (Name[0] ==
'a' && Name.back() ==
'4')
5029 IID = Intrinsic::x86_addcarry_64;
5030 else if (Name[0] ==
's' && Name.back() ==
'2')
5031 IID = Intrinsic::x86_subborrow_32;
5032 else if (Name[0] ==
's' && Name.back() ==
'4')
5033 IID = Intrinsic::x86_subborrow_64;
5040 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5043 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5046 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5050 }
else if (Name.starts_with(
"avx512.mask.") &&
5053 }
else if (Name.starts_with(
"bmi.pdep.")) {
5055 }
else if (Name.starts_with(
"bmi.pext.")) {
5065 if (Name.starts_with(
"neon.bfcvt")) {
5066 if (Name.starts_with(
"neon.bfcvtn2")) {
5068 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5070 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5071 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5074 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5075 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5077 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5081 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5082 return Builder.CreateShuffleVector(
5085 return Builder.CreateFPTrunc(CI->
getOperand(0),
5088 }
else if (Name.starts_with(
"sve.fcvt")) {
5091 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5092 .
Case(
"sve.fcvtnt.bf16f32",
5093 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5105 if (Args[1]->
getType() != BadPredTy)
5108 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5109 BadPredTy, Args[1]);
5110 Args[1] = Builder.CreateIntrinsic(
5111 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5113 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5117 if (Name ==
"neon.vcvtfp2hf")
5118 return Builder.CreateBitCast(
5119 Builder.CreateFPTrunc(
5123 if (Name ==
"neon.vcvthf2fp")
5124 return Builder.CreateFPExt(
5125 Builder.CreateBitCast(
5135 if (Name ==
"mve.vctp64.old") {
5138 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5141 Value *C1 = Builder.CreateIntrinsic(
5142 Intrinsic::arm_mve_pred_v2i,
5144 return Builder.CreateIntrinsic(
5145 Intrinsic::arm_mve_pred_i2v,
5147 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5148 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5149 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5150 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5152 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5153 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5154 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5155 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5157 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5158 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5159 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5160 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5161 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5162 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5163 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5164 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5165 std::vector<Type *> Tys;
5169 case Intrinsic::arm_mve_mull_int_predicated:
5170 case Intrinsic::arm_mve_vqdmull_predicated:
5171 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5174 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5175 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5176 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5180 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5184 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5188 case Intrinsic::arm_cde_vcx1q_predicated:
5189 case Intrinsic::arm_cde_vcx1qa_predicated:
5190 case Intrinsic::arm_cde_vcx2q_predicated:
5191 case Intrinsic::arm_cde_vcx2qa_predicated:
5192 case Intrinsic::arm_cde_vcx3q_predicated:
5193 case Intrinsic::arm_cde_vcx3qa_predicated:
5200 std::vector<Value *>
Ops;
5202 Type *Ty =
Op->getType();
5203 if (Ty->getScalarSizeInBits() == 1) {
5204 Value *C1 = Builder.CreateIntrinsic(
5205 Intrinsic::arm_mve_pred_v2i,
5207 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5212 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5227 auto UpgradeLegacyWMMAIUIntrinsicCall =
5232 Args.push_back(Builder.getFalse());
5236 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5243 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5247 NewCall->copyMetadata(*CI);
5251 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5252 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5253 "intrinsic should have 7 arguments");
5256 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5258 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5259 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5260 "intrinsic should have 8 arguments");
5265 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5268 switch (
F->getIntrinsicID()) {
5271 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5272 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5273 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5274 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5275 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5276 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5291 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5294 F->getParent(),
F->getIntrinsicID(), Overloads);
5299 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5303 NewCall->copyMetadata(*CI);
5304 NewCall->takeName(CI);
5309 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5315 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5316 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5324 if (Name.starts_with(
"addrspacecast.nonnull")) {
5327 Value *ASC = Builder.CreateAddrSpaceCast(
5350 if (NumOperands < 3)
5363 bool IsVolatile =
false;
5367 if (NumOperands > 3)
5372 if (NumOperands > 5) {
5374 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5388 if (VT->getElementType()->isIntegerTy(16)) {
5391 Val = Builder.CreateBitCast(Val, AsBF16);
5399 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5401 unsigned AddrSpace = PtrTy->getAddressSpace();
5404 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5406 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5411 MDNode *RangeNotPrivate =
5414 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5420 return Builder.CreateBitCast(RMW, RetTy);
5441 return MAV->getMetadata();
5450 if (Name ==
"label") {
5452 }
else if (Name ==
"assign") {
5459 }
else if (Name ==
"declare") {
5463 }
else if (Name ==
"addr") {
5473 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5474 }
else if (Name ==
"value") {
5477 unsigned ExprOp = 2;
5492 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5500 int64_t OffsetVal =
Offset->getSExtValue();
5501 return Builder.CreateIntrinsic(OffsetVal >= 0
5502 ? Intrinsic::vector_splice_left
5503 : Intrinsic::vector_splice_right,
5505 {CI->getArgOperand(0), CI->getArgOperand(1),
5506 Builder.getInt32(std::abs(OffsetVal))});
5511 if (Name.starts_with(
"to.fp16")) {
5513 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5514 return Builder.CreateBitCast(Cast, CI->
getType());
5517 if (Name.starts_with(
"from.fp16")) {
5519 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5520 return Builder.CreateFPExt(Cast, CI->
getType());
5579 else if (Opcode == Instruction::ICmp)
5582 else if (Opcode == Instruction::FCmp)
5585 else if (Opcode == Instruction::Select)
5590 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5602 if (Defaults.empty())
5605 unsigned OldArgCount = CI->
arg_size();
5606 unsigned NewArgCount = NewFn->
arg_size();
5610 if (OldArgCount >= NewArgCount)
5618 if (OldArgCount < FirstDefault)
5623 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5624 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5625 "missing argument outside the default range");
5626 Type *ParamTy = NewFT->getParamType(Idx);
5631 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5637 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5669 if (!Name.consume_front(
"llvm."))
5672 bool IsX86 = Name.consume_front(
"x86.");
5673 bool IsNVVM = Name.consume_front(
"nvvm.");
5674 bool IsAArch64 = Name.consume_front(
"aarch64.");
5675 bool IsARM = Name.consume_front(
"arm.");
5676 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5677 bool IsDbg = Name.consume_front(
"dbg.");
5679 (Name.consume_front(
"experimental.vector.splice") ||
5680 Name.consume_front(
"vector.splice")) &&
5681 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5682 Value *Rep =
nullptr;
5684 if (!IsX86 && Name ==
"stackprotectorcheck") {
5686 }
else if (IsNVVM) {
5690 }
else if (IsAArch64) {
5694 }
else if (IsAMDGCN) {
5698 }
else if (IsOldSplice) {
5700 }
else if (Name.consume_front(
"convert.")) {
5702 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5717 const auto &DefaultCase = [&]() ->
void {
5725 "Unknown function for CallBase upgrade and isn't just a name change");
5733 "Return type must have changed");
5734 assert(OldST->getNumElements() ==
5736 "Must have same number of elements");
5739 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5742 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5743 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5744 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5768 case Intrinsic::arm_neon_vst1:
5769 case Intrinsic::arm_neon_vst2:
5770 case Intrinsic::arm_neon_vst3:
5771 case Intrinsic::arm_neon_vst4:
5772 case Intrinsic::arm_neon_vst2lane:
5773 case Intrinsic::arm_neon_vst3lane:
5774 case Intrinsic::arm_neon_vst4lane: {
5776 NewCall = Builder.CreateCall(NewFn, Args);
5779 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5780 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5781 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5786 NewCall = Builder.CreateCall(NewFn, Args);
5789 case Intrinsic::aarch64_sve_ld3_sret:
5790 case Intrinsic::aarch64_sve_ld4_sret:
5791 case Intrinsic::aarch64_sve_ld2_sret: {
5799 Name = Name.substr(5);
5806 unsigned MinElts = RetTy->getMinNumElements() /
N;
5808 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5810 for (
unsigned I = 0;
I <
N;
I++) {
5811 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5812 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5818 case Intrinsic::coro_end_async:
5819 case Intrinsic::coro_end: {
5821 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5823 NewCall = Builder.CreateCall(NewFn, Args);
5828 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5829 Value *InRamp = Builder.CreateCall(IsInRamp);
5839 case Intrinsic::vector_extract: {
5841 Name = Name.substr(5);
5842 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5847 unsigned MinElts = RetTy->getMinNumElements();
5850 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5854 case Intrinsic::vector_insert: {
5856 Name = Name.substr(5);
5857 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5861 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5866 NewCall = Builder.CreateCall(
5870 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5876 assert(
N > 1 &&
"Create is expected to be between 2-4");
5879 unsigned MinElts = RetTy->getMinNumElements() /
N;
5880 for (
unsigned I = 0;
I <
N;
I++) {
5882 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5889 case Intrinsic::arm_neon_bfdot:
5890 case Intrinsic::arm_neon_bfmmla:
5891 case Intrinsic::arm_neon_bfmlalb:
5892 case Intrinsic::arm_neon_bfmlalt:
5893 case Intrinsic::aarch64_neon_bfdot:
5894 case Intrinsic::aarch64_neon_bfmmla:
5895 case Intrinsic::aarch64_neon_bfmlalb:
5896 case Intrinsic::aarch64_neon_bfmlalt: {
5899 "Mismatch between function args and call args");
5900 size_t OperandWidth =
5902 assert((OperandWidth == 64 || OperandWidth == 128) &&
5903 "Unexpected operand width");
5905 auto Iter = CI->
args().begin();
5906 Args.push_back(*Iter++);
5907 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5908 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5909 NewCall = Builder.CreateCall(NewFn, Args);
5913 case Intrinsic::bitreverse:
5914 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5917 case Intrinsic::ctlz:
5918 case Intrinsic::cttz: {
5925 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5929 case Intrinsic::objectsize: {
5930 Value *NullIsUnknownSize =
5934 NewCall = Builder.CreateCall(
5939 case Intrinsic::ctpop:
5940 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5942 case Intrinsic::dbg_value: {
5944 Name = Name.substr(5);
5946 if (Name.starts_with(
"dbg.addr")) {
5960 if (
Offset->isNullValue()) {
5961 NewCall = Builder.CreateCall(
5970 case Intrinsic::ptr_annotation:
5978 NewCall = Builder.CreateCall(
5987 case Intrinsic::var_annotation:
5994 NewCall = Builder.CreateCall(
6003 case Intrinsic::riscv_aes32dsi:
6004 case Intrinsic::riscv_aes32dsmi:
6005 case Intrinsic::riscv_aes32esi:
6006 case Intrinsic::riscv_aes32esmi:
6007 case Intrinsic::riscv_sm4ks:
6008 case Intrinsic::riscv_sm4ed: {
6018 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6019 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6025 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6026 Value *Res = NewCall;
6028 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6034 case Intrinsic::nvvm_mapa_shared_cluster: {
6038 Value *Res = NewCall;
6039 Res = Builder.CreateAddrSpaceCast(
6046 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
6047 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6050 Args[0] = Builder.CreateAddrSpaceCast(
6053 NewCall = Builder.CreateCall(NewFn, Args);
6060#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6061 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6063#undef G2S_CLUSTER_CASE
6068 Args[0] = Builder.CreateAddrSpaceCast(
6074 Args.push_back(Builder.getInt32(0));
6076 NewCall = Builder.CreateCall(NewFn, Args);
6083#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6084 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6092 "expected only the trailing validate_pattern to be missing");
6093 Args.push_back(Builder.getInt32(0));
6095 NewCall = Builder.CreateCall(NewFn, Args);
6101#undef NVVM_TMA_G2S_MODES
6104 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6105 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6106 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6107 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6108 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6109 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6110 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6111 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6113 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6117 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6118 NewCall = Builder.CreateCall(NewFn, Args);
6121 case Intrinsic::nvvm_tcgen05_mma_shared:
6122 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
6123 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
6124 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
6125 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
6126 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
6127 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
6128 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
6129 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
6130 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
6131 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
6132 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
6133 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
6134 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
6135 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
6136 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
6137 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
6138 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
6139 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
6140 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
6141 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
6142 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
6143 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
6144 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
6145 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
6146 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
6147 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
6148 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
6149 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
6150 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
6151 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
6152 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
6153 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
6154 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
6155 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
6156 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
6157 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
6158 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
6159 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
6161 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
6162 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
6164 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
6165 case Intrinsic::nvvm_tcgen05_mma_tensor:
6166 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
6167 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
6168 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
6169 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
6170 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
6171 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
6172 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
6173 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
6174 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
6175 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
6176 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
6177 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
6178 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
6179 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
6181 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
6182 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
6184 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
6186 Args.push_back(Builder.getInt32(0));
6187 NewCall = Builder.CreateCall(NewFn, Args);
6190 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6191 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6192 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6193 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6196 Builder.getFalse()});
6198 case Intrinsic::riscv_sha256sig0:
6199 case Intrinsic::riscv_sha256sig1:
6200 case Intrinsic::riscv_sha256sum0:
6201 case Intrinsic::riscv_sha256sum1:
6202 case Intrinsic::riscv_sm3p0:
6203 case Intrinsic::riscv_sm3p1: {
6210 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6212 NewCall = Builder.CreateCall(NewFn, Arg);
6214 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6221 case Intrinsic::x86_xop_vfrcz_ss:
6222 case Intrinsic::x86_xop_vfrcz_sd:
6223 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6226 case Intrinsic::x86_xop_vpermil2pd:
6227 case Intrinsic::x86_xop_vpermil2ps:
6228 case Intrinsic::x86_xop_vpermil2pd_256:
6229 case Intrinsic::x86_xop_vpermil2ps_256: {
6233 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6234 NewCall = Builder.CreateCall(NewFn, Args);
6238 case Intrinsic::x86_sse41_ptestc:
6239 case Intrinsic::x86_sse41_ptestz:
6240 case Intrinsic::x86_sse41_ptestnzc: {
6254 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6255 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6257 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6261 case Intrinsic::x86_rdtscp: {
6267 NewCall = Builder.CreateCall(NewFn);
6269 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6272 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6280 case Intrinsic::x86_sse41_insertps:
6281 case Intrinsic::x86_sse41_dppd:
6282 case Intrinsic::x86_sse41_dpps:
6283 case Intrinsic::x86_sse41_mpsadbw:
6284 case Intrinsic::x86_avx_dp_ps_256:
6285 case Intrinsic::x86_avx2_mpsadbw: {
6291 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6292 NewCall = Builder.CreateCall(NewFn, Args);
6296 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6297 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6298 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6299 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6300 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6301 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6307 NewCall = Builder.CreateCall(NewFn, Args);
6316 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6317 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6318 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6319 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6320 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6321 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6325 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6326 Args[1] = Builder.CreateBitCast(
6329 NewCall = Builder.CreateCall(NewFn, Args);
6330 Value *Res = Builder.CreateBitCast(
6338 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6339 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6340 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6344 Args[1] = Builder.CreateBitCast(
6346 Args[2] = Builder.CreateBitCast(
6349 NewCall = Builder.CreateCall(NewFn, Args);
6353 case Intrinsic::thread_pointer: {
6354 NewCall = Builder.CreateCall(NewFn, {});
6358 case Intrinsic::memcpy:
6359 case Intrinsic::memmove:
6360 case Intrinsic::memset: {
6376 NewCall = Builder.CreateCall(NewFn, Args);
6378 AttributeList NewAttrs = AttributeList::get(
6379 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6380 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6381 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6386 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6389 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6393 case Intrinsic::masked_load:
6394 case Intrinsic::masked_gather:
6395 case Intrinsic::masked_store:
6396 case Intrinsic::masked_scatter: {
6402 auto GetMaybeAlign = [](
Value *
Op) {
6404 uint64_t Val = CI->getZExtValue();
6412 auto GetAlign = [&](
Value *
Op) {
6421 case Intrinsic::masked_load:
6422 NewCall = Builder.CreateMaskedLoad(
6426 case Intrinsic::masked_gather:
6427 NewCall = Builder.CreateMaskedGather(
6433 case Intrinsic::masked_store:
6434 NewCall = Builder.CreateMaskedStore(
6438 case Intrinsic::masked_scatter:
6439 NewCall = Builder.CreateMaskedScatter(
6441 DL.getValueOrABITypeAlignment(
6455 case Intrinsic::lifetime_start:
6456 case Intrinsic::lifetime_end: {
6468 NewCall = Builder.CreateLifetimeStart(Ptr);
6470 NewCall = Builder.CreateLifetimeEnd(Ptr);
6479 case Intrinsic::x86_avx512_vpdpbusd_128:
6480 case Intrinsic::x86_avx512_vpdpbusd_256:
6481 case Intrinsic::x86_avx512_vpdpbusd_512:
6482 case Intrinsic::x86_avx512_vpdpbusds_128:
6483 case Intrinsic::x86_avx512_vpdpbusds_256:
6484 case Intrinsic::x86_avx512_vpdpbusds_512:
6485 case Intrinsic::x86_avx2_vpdpbssd_128:
6486 case Intrinsic::x86_avx2_vpdpbssd_256:
6487 case Intrinsic::x86_avx10_vpdpbssd_512:
6488 case Intrinsic::x86_avx2_vpdpbssds_128:
6489 case Intrinsic::x86_avx2_vpdpbssds_256:
6490 case Intrinsic::x86_avx10_vpdpbssds_512:
6491 case Intrinsic::x86_avx2_vpdpbsud_128:
6492 case Intrinsic::x86_avx2_vpdpbsud_256:
6493 case Intrinsic::x86_avx10_vpdpbsud_512:
6494 case Intrinsic::x86_avx2_vpdpbsuds_128:
6495 case Intrinsic::x86_avx2_vpdpbsuds_256:
6496 case Intrinsic::x86_avx10_vpdpbsuds_512:
6497 case Intrinsic::x86_avx2_vpdpbuud_128:
6498 case Intrinsic::x86_avx2_vpdpbuud_256:
6499 case Intrinsic::x86_avx10_vpdpbuud_512:
6500 case Intrinsic::x86_avx2_vpdpbuuds_128:
6501 case Intrinsic::x86_avx2_vpdpbuuds_256:
6502 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6507 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6508 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6510 NewCall = Builder.CreateCall(NewFn, Args);
6513 case Intrinsic::x86_avx512_vpdpwssd_128:
6514 case Intrinsic::x86_avx512_vpdpwssd_256:
6515 case Intrinsic::x86_avx512_vpdpwssd_512:
6516 case Intrinsic::x86_avx512_vpdpwssds_128:
6517 case Intrinsic::x86_avx512_vpdpwssds_256:
6518 case Intrinsic::x86_avx512_vpdpwssds_512:
6519 case Intrinsic::x86_avx2_vpdpwsud_128:
6520 case Intrinsic::x86_avx2_vpdpwsud_256:
6521 case Intrinsic::x86_avx10_vpdpwsud_512:
6522 case Intrinsic::x86_avx2_vpdpwsuds_128:
6523 case Intrinsic::x86_avx2_vpdpwsuds_256:
6524 case Intrinsic::x86_avx10_vpdpwsuds_512:
6525 case Intrinsic::x86_avx2_vpdpwusd_128:
6526 case Intrinsic::x86_avx2_vpdpwusd_256:
6527 case Intrinsic::x86_avx10_vpdpwusd_512:
6528 case Intrinsic::x86_avx2_vpdpwusds_128:
6529 case Intrinsic::x86_avx2_vpdpwusds_256:
6530 case Intrinsic::x86_avx10_vpdpwusds_512:
6531 case Intrinsic::x86_avx2_vpdpwuud_128:
6532 case Intrinsic::x86_avx2_vpdpwuud_256:
6533 case Intrinsic::x86_avx10_vpdpwuud_512:
6534 case Intrinsic::x86_avx2_vpdpwuuds_128:
6535 case Intrinsic::x86_avx2_vpdpwuuds_256:
6536 case Intrinsic::x86_avx10_vpdpwuuds_512:
6541 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6542 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6544 NewCall = Builder.CreateCall(NewFn, Args);
6547 assert(NewCall &&
"Should have either set this variable or returned through "
6548 "the default case");
6555 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6569 F->eraseFromParent();
6575 if (NumOperands == 0)
6583 if (NumOperands == 3) {
6587 Metadata *Elts2[] = {ScalarType, ScalarType,
6601 if (
Opc != Instruction::BitCast)
6605 Type *SrcTy = V->getType();
6622 if (
Opc != Instruction::BitCast)
6625 Type *SrcTy =
C->getType();
6642 if (Flag.getNumOperands() < 3)
6643 return std::nullopt;
6645 return Name->getString();
6646 return std::nullopt;
6660 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6661 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6662 if (auto Name = getModuleFlagNameSafely(*Flag))
6663 return *Name ==
"Debug Info Version";
6666 if (OpIt != ModFlags->op_end()) {
6667 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6674 bool BrokenDebugInfo =
false;
6677 if (!BrokenDebugInfo)
6683 M.getContext().diagnose(Diag);
6690 M.getContext().diagnose(DiagVersion);
6700 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6703 if (
F->hasFnAttribute(Attr)) {
6706 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6708 auto [Part, Rest] = S.
split(
',');
6714 const unsigned Dim = DimC -
'x';
6715 assert(Dim < 3 &&
"Unexpected dim char");
6725 F->addFnAttr(Attr, NewAttr);
6729 return S ==
"x" || S ==
"y" || S ==
"z";
6734 if (K ==
"kernel") {
6746 const unsigned Idx = (AlignIdxValuePair >> 16);
6747 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6752 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6757 if (K ==
"minctasm") {
6762 if (K ==
"maxnreg") {
6767 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6771 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6775 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6779 if (K ==
"grid_constant") {
6794 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6801 if (!SeenNodes.
insert(MD).second)
6808 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6815 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6817 const MDOperand &V = MD->getOperand(j + 1);
6820 NewOperands.
append({K, V});
6823 if (NewOperands.
size() > 1)
6836 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6837 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6838 if (ModRetainReleaseMarker) {
6844 ID->getString().split(ValueComp,
"#");
6845 if (ValueComp.
size() == 2) {
6846 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6850 M.eraseNamedMetadata(ModRetainReleaseMarker);
6861 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6887 bool InvalidCast =
false;
6889 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6902 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6904 Args.push_back(Arg);
6911 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6916 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6929 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6937 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6938 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6939 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6940 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6941 {
"objc_autoreleaseReturnValue",
6942 llvm::Intrinsic::objc_autoreleaseReturnValue},
6943 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6944 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6945 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6946 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6947 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6948 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6949 {
"objc_release", llvm::Intrinsic::objc_release},
6950 {
"objc_retain", llvm::Intrinsic::objc_retain},
6951 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6952 {
"objc_retainAutoreleaseReturnValue",
6953 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6954 {
"objc_retainAutoreleasedReturnValue",
6955 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6956 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6957 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6958 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6959 {
"objc_unsafeClaimAutoreleasedReturnValue",
6960 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6961 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6962 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6963 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6964 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6965 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6966 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6967 {
"objc_arc_annotation_topdown_bbstart",
6968 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6969 {
"objc_arc_annotation_topdown_bbend",
6970 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6971 {
"objc_arc_annotation_bottomup_bbstart",
6972 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6973 {
"objc_arc_annotation_bottomup_bbend",
6974 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6976 for (
auto &
I : RuntimeFuncs)
6977 UpgradeToIntrinsic(
I.first,
I.second);
7001 std::optional<bool> UseAddressDisc;
7004 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7005 for (
const MDNode *Flag : ModFlags->operands()) {
7007 if (Name && (*Name ==
"ptrauth-init-fini" ||
7008 *Name ==
"ptrauth-init-fini-address-discrimination"))
7013 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7014 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7015 constexpr unsigned ExpectedAddressMarker = 1;
7018 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7021 bool HasAddressDisc;
7022 if (!CPA->hasAddressDiscriminator())
7023 HasAddressDisc =
false;
7024 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7025 HasAddressDisc =
true;
7029 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7032 UseAddressDisc = HasAddressDisc;
7033 return CPA->getPointer();
7037 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7040 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7042 if (!GV || !GV->hasInitializer())
7046 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7049 std::vector<Constant *> NewStructors;
7050 NewStructors.reserve(OldStructorsArray->getNumOperands());
7052 for (
Use &U : OldStructorsArray->operands()) {
7061 Func = UpgradeSinglePointer(Func);
7065 NewStructors.push_back(
7074 if (GlobalArraysToUpgrade.
empty())
7076 assert(UseAddressDisc.has_value());
7078 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7079 GV->setInitializer(NewInit);
7082 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7092 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7096 bool HasObjCFlag =
false, HasClassProperties =
false;
7097 bool HasSwiftVersionFlag =
false;
7098 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7105 if (
Op->getNumOperands() != 3)
7119 if (ID->getString() ==
"Objective-C Image Info Version")
7121 if (ID->getString() ==
"Objective-C Class Properties")
7122 HasClassProperties =
true;
7124 if (ID->getString() ==
"PIC Level") {
7125 if (
auto *Behavior =
7127 uint64_t V = Behavior->getLimitedValue();
7133 if (ID->getString() ==
"PIE Level")
7134 if (
auto *Behavior =
7141 if (ID->getString() ==
"branch-target-enforcement" ||
7142 ID->getString().starts_with(
"sign-return-address")) {
7143 if (
auto *Behavior =
7149 Op->getOperand(1),
Op->getOperand(2)};
7159 if (ID->getString() ==
"Objective-C Image Info Section") {
7162 Value->getString().split(ValueComp,
" ");
7163 if (ValueComp.
size() != 1) {
7164 std::string NewValue;
7165 for (
auto &S : ValueComp)
7166 NewValue += S.str();
7177 if (ID->getString() ==
"Objective-C Garbage Collection") {
7180 assert(Md->getValue() &&
"Expected non-empty metadata");
7181 auto Type = Md->getValue()->getType();
7184 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7185 if ((Val & 0xff) != Val) {
7186 HasSwiftVersionFlag =
true;
7187 SwiftABIVersion = (Val & 0xff00) >> 8;
7188 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7189 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7200 if (ID->getString() ==
"amdgpu_code_object_version") {
7203 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7212 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7241 if (HasObjCFlag && !HasClassProperties) {
7247 if (HasSwiftVersionFlag) {
7251 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7253 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7261 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7265 auto MatchesVersion = [](
const MDNode *
Op) {
7266 return Op->getNumOperands() >= 3 &&
7280 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7281 assert(
Op->getNumOperands() >= 2 &&
7282 "Expected at least 2 operands - name and linkage type");
7294 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7305 auto TrimSpaces = [](
StringRef Section) -> std::string {
7307 Section.split(Components,
',');
7312 for (
auto Component : Components)
7313 OS <<
',' << Component.trim();
7318 for (
auto &GV : M.globals()) {
7319 if (!GV.hasSection())
7324 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7329 GV.setSection(TrimSpaces(Section));
7345struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7346 StrictFPUpgradeVisitor() =
default;
7349 if (!
Call.isStrictFP())
7355 Call.removeFnAttr(Attribute::StrictFP);
7356 Call.addFnAttr(Attribute::NoBuiltin);
7361struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7362 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7363 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7365 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7380 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7381 StrictFPUpgradeVisitor SFPV;
7386 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7387 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7388 for (
auto &Arg :
F.args())
7390 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7392 bool AddingAttrs =
false, RemovingAttrs =
false;
7393 AttrBuilder AttrsToAdd(
F.getContext());
7398 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7399 A.isValid() &&
A.isStringAttribute()) {
7400 F.setSection(
A.getValueAsString());
7402 RemovingAttrs =
true;
7406 A.isValid() &&
A.isStringAttribute()) {
7409 AddingAttrs = RemovingAttrs =
true;
7412 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7413 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7415 RemovingAttrs =
true;
7416 if (
A.getValueAsString() ==
"true") {
7417 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7426 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7429 if (
A.getValueAsBool()) {
7430 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7436 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7437 RemovingAttrs =
true;
7444 bool HandleDenormalMode =
false;
7446 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7449 DenormalFPMath = ParsedMode;
7451 AddingAttrs = RemovingAttrs =
true;
7452 HandleDenormalMode =
true;
7456 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7460 DenormalFPMathF32 = ParsedMode;
7462 AddingAttrs = RemovingAttrs =
true;
7463 HandleDenormalMode =
true;
7467 if (HandleDenormalMode)
7468 AttrsToAdd.addDenormalFPEnvAttr(
7472 F.removeFnAttrs(AttrsToRemove);
7475 F.addFnAttrs(AttrsToAdd);
7481 if (!
F.hasFnAttribute(FnAttrName))
7482 F.addFnAttr(FnAttrName,
Value);
7489 if (!
F.hasFnAttribute(FnAttrName)) {
7491 F.addFnAttr(FnAttrName);
7493 auto A =
F.getFnAttribute(FnAttrName);
7494 if (
"false" ==
A.getValueAsString())
7495 F.removeFnAttr(FnAttrName);
7496 else if (
"true" ==
A.getValueAsString()) {
7497 F.removeFnAttr(FnAttrName);
7498 F.addFnAttr(FnAttrName);
7504 Triple T(M.getTargetTriple());
7505 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7508 uint64_t BTEValue = 0;
7509 uint64_t BPPLRValue = 0;
7510 uint64_t GCSValue = 0;
7511 uint64_t SRAValue = 0;
7512 uint64_t SRAALLValue = 0;
7513 uint64_t SRABKeyValue = 0;
7515 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7519 if (
Op->getNumOperands() != 3)
7528 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7529 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7530 : IDStr ==
"guarded-control-stack" ? &GCSValue
7531 : IDStr ==
"sign-return-address" ? &SRAValue
7532 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7533 : IDStr ==
"sign-return-address-with-bkey"
7539 *ValPtr = CI->getZExtValue();
7545 bool BTE = BTEValue == 1;
7546 bool BPPLR = BPPLRValue == 1;
7547 bool GCS = GCSValue == 1;
7548 bool SRA = SRAValue == 1;
7551 if (SRA && SRAALLValue == 1)
7552 SignTypeValue =
"all";
7555 if (SRA && SRABKeyValue == 1)
7556 SignKeyValue =
"b_key";
7558 for (
Function &
F : M.getFunctionList()) {
7559 if (
F.isDeclaration())
7566 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7567 A.isValid() &&
"none" ==
A.getValueAsString()) {
7568 F.removeFnAttr(
"sign-return-address");
7569 F.removeFnAttr(
"sign-return-address-key");
7585 if (SRAALLValue == 1)
7587 if (SRABKeyValue == 1)
7614 if (
T->getNumOperands() < 1)
7619 if (S->getString().starts_with(
"llvm.vectorizer."))
7625 StringRef OldPrefix =
"llvm.vectorizer.";
7628 if (OldTag ==
"llvm.vectorizer.unroll")
7640 if (
T->getNumOperands() < 1)
7652 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7665 Ops.reserve(
T->getNumOperands());
7666 Ops.push_back(NewTag);
7667 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7668 Ops.push_back(
T->getOperand(
I));
7685 if (
T->isDistinct()) {
7686 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7698 Ops.reserve(
T->getNumOperands());
7709 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7710 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7711 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7714 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7716 auto I =
DL.find(
"-n64-");
7718 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7723 std::string Res =
DL.str();
7726 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7727 Res.append(Res.empty() ?
"G1" :
"-G1");
7735 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7736 Res.append(
"-ni:7:8:9");
7738 if (
DL.ends_with(
"ni:7"))
7740 if (
DL.ends_with(
"ni:7:8"))
7745 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7746 Res.append(
"-p7:160:256:256:32");
7747 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7748 Res.append(
"-p8:128:128:128:48");
7749 constexpr StringRef OldP8(
"-p8:128:128-");
7750 if (
DL.contains(OldP8))
7751 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7752 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7753 Res.append(
"-p9:192:256:256:32");
7757 if (!
DL.contains(
"m:e"))
7758 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7763 if (
T.isSystemZ() && !
DL.empty()) {
7765 if (!
DL.contains(
"-S64"))
7766 return "E-S64" +
DL.drop_front(1).str();
7770 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7773 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7774 if (!
DL.contains(AddrSpaces)) {
7776 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7777 if (R.match(Res, &
Groups))
7783 if (
T.isAArch64()) {
7785 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7786 Res.append(
"-Fn32");
7787 AddPtr32Ptr64AddrSpaces();
7791 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7795 std::string I64 =
"-i64:64";
7796 std::string I128 =
"-i128:128";
7798 size_t Pos = Res.find(I64);
7799 if (Pos !=
size_t(-1))
7800 Res.insert(Pos + I64.size(), I128);
7804 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7805 size_t Pos = Res.find(
"-S128");
7808 Res.insert(Pos,
"-f64:32:64");
7814 AddPtr32Ptr64AddrSpaces();
7822 if (!
T.isOSIAMCU()) {
7823 std::string I128 =
"-i128:128";
7826 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7827 if (R.match(Res, &
Groups))
7835 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7837 auto I =
Ref.find(
"-f80:32-");
7839 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7847 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7850 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7851 B.removeAttribute(
"no-frame-pointer-elim");
7853 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7855 if (FramePointer !=
"all")
7856 FramePointer =
"non-leaf";
7857 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7859 if (!FramePointer.
empty())
7860 B.addAttribute(
"frame-pointer", FramePointer);
7862 A =
B.getAttribute(
"null-pointer-is-valid");
7865 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7866 B.removeAttribute(
"null-pointer-is-valid");
7867 if (NullPointerIsValid)
7868 B.addAttribute(Attribute::NullPointerIsValid);
7871 A =
B.getAttribute(
"uniform-work-group-size");
7875 bool IsTrue = Val ==
"true";
7876 B.removeAttribute(
"uniform-work-group-size");
7878 B.addAttribute(
"uniform-work-group-size");
7889 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#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 Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05MMAIntrinsic(Function *F, StringRef Name)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#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 Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static 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 ID lookupIntrinsicID(StringRef Name)
This does the actual lookup of an intrinsic ID which matches the given function name.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
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.