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 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1505 return Intrinsic::nvvm_mbarrier_init;
1509 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1510 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1511 Name.consume_front(
"param");
1515 if (!Name.consume_front(
"vp."))
1544 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1545 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1552 if (!Name.consume_front(
"vp."))
1572 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1573 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1578 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1590 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1601 if (Name.starts_with(
"to.fp16")) {
1605 FuncTy->getReturnType());
1608 if (Name.starts_with(
"from.fp16")) {
1612 FuncTy->getReturnType());
1621 if (Defaults.empty())
1627 unsigned FullArgCount = FirstDefault + Defaults.size();
1630 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1633 return FullArgCount;
1640 if (FullArgCount == 0)
1646 "total number of default args does not match intrinsic signature");
1651 bool CanUpgradeDebugIntrinsicsToRecords) {
1652 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1657 if (!Name.consume_front(
"llvm.") || Name.empty())
1663 bool IsArm = Name.consume_front(
"arm.");
1664 if (IsArm || Name.consume_front(
"aarch64.")) {
1670 if (Name.consume_front(
"amdgcn.")) {
1671 if (Name ==
"alignbit") {
1674 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1678 if (Name.consume_front(
"atomic.")) {
1679 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1680 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1689 if (Name.starts_with(
"addrspacecast.nonnull")) {
1696 switch (
F->getIntrinsicID()) {
1700 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1701 if (
F->arg_size() == 7) {
1706 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1707 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1708 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1709 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1710 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1711 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1712 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1713 if (
F->arg_size() == 8) {
1720 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1721 Name.consume_front(
"flat.atomic.")) {
1722 if (Name.starts_with(
"fadd") ||
1724 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1725 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1733 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1738 if (Name.starts_with(
"ldexp.")) {
1741 F->getParent(), Intrinsic::ldexp,
1742 {F->getReturnType(), F->getArg(1)->getType()});
1751 if (
F->arg_size() == 1) {
1752 if (Name.consume_front(
"convert.")) {
1766 F->arg_begin()->getType());
1772 if (Name ==
"coro.end" &&
1773 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1774 CoroEndID = Intrinsic::coro_end;
1775 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1776 CoroEndID = Intrinsic::coro_end_async;
1787 if (Name.consume_front(
"dbg.")) {
1789 if (CanUpgradeDebugIntrinsicsToRecords) {
1790 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1791 Name ==
"declare" || Name ==
"label") {
1800 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1803 Intrinsic::dbg_value);
1810 if (Name.consume_front(
"experimental.vector.")) {
1816 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1817 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1818 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1819 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1820 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1822 Intrinsic::vector_partial_reduce_add)
1825 const auto *FT =
F->getFunctionType();
1827 if (ID == Intrinsic::vector_extract ||
1828 ID == Intrinsic::vector_interleave2)
1831 if (ID != Intrinsic::vector_interleave2)
1833 if (ID == Intrinsic::vector_insert ||
1834 ID == Intrinsic::vector_partial_reduce_add)
1842 if (Name.consume_front(
"reduce.")) {
1844 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1845 if (R.match(Name, &
Groups))
1847 .
Case(
"add", Intrinsic::vector_reduce_add)
1848 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1849 .
Case(
"and", Intrinsic::vector_reduce_and)
1850 .
Case(
"or", Intrinsic::vector_reduce_or)
1851 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1852 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1853 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1854 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1855 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1856 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1857 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1862 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1867 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1868 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1873 auto Args =
F->getFunctionType()->params();
1875 {Args[V2 ? 1 : 0]});
1881 if (Name.consume_front(
"splice"))
1885 if (Name.consume_front(
"experimental.stepvector.")) {
1889 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1894 if (Name.starts_with(
"flt.rounds")) {
1897 Intrinsic::get_rounding);
1902 if (Name.starts_with(
"invariant.group.barrier")) {
1904 auto Args =
F->getFunctionType()->params();
1905 Type* ObjectPtr[1] = {Args[0]};
1908 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1913 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1914 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1915 if (IsLifetimeStart || IsLifetimeEnd) {
1916 if (
F->arg_size() == 2) {
1917 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1918 : Intrinsic::lifetime_end;
1923 F->getArg(1)->getType());
1925 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1945 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1946 .StartsWith(
"memmove.", Intrinsic::memmove)
1948 if (
F->arg_size() == 5) {
1952 F->getFunctionType()->params().slice(0, 3);
1958 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1961 const auto *FT =
F->getFunctionType();
1962 Type *ParamTypes[2] = {
1963 FT->getParamType(0),
1967 Intrinsic::memset, ParamTypes);
1973 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1974 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1975 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1976 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1978 if (MaskedID &&
F->arg_size() == 4) {
1980 if (MaskedID == Intrinsic::masked_load ||
1981 MaskedID == Intrinsic::masked_gather) {
1983 F->getParent(), MaskedID,
1984 {F->getReturnType(), F->getArg(0)->getType()});
1988 F->getParent(), MaskedID,
1989 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1995 if (Name.consume_front(
"nvvm.")) {
1997 if (
F->arg_size() == 1) {
2000 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2001 .Case(
"clz.i", Intrinsic::ctlz)
2002 .
Case(
"popc.i", Intrinsic::ctpop)
2006 {F->getReturnType()});
2009 }
else if (
F->arg_size() == 2) {
2012 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2013 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2014 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2015 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2019 {F->getReturnType()});
2056 F->getParent(), IID,
F->getReturnType(),
2057 F->getFunctionType()->params());
2068 {F->getArg(0)->getType()});
2103 F->getArg(0)->getType());
2111 bool Expand =
false;
2112 if (Name.consume_front(
"abs."))
2115 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2116 else if (Name.consume_front(
"fabs."))
2118 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2119 else if (Name.consume_front(
"add."))
2122 else if (Name.consume_front(
"ex2.approx."))
2125 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2126 else if (Name.consume_front(
"atomic.load."))
2135 else if (Name.consume_front(
"atomic."))
2150 else if (Name.consume_front(
"bitcast."))
2153 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2154 else if (Name.consume_front(
"rotate."))
2156 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2157 else if (Name.consume_front(
"ptr.gen.to."))
2160 else if (Name.consume_front(
"ptr."))
2163 else if (Name.consume_front(
"ldg.global."))
2165 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2166 Name.starts_with(
"p."));
2169 .
Case(
"barrier0",
true)
2170 .
Case(
"barrier.n",
true)
2171 .
Case(
"barrier.sync.cnt",
true)
2172 .
Case(
"barrier.sync",
true)
2173 .
Case(
"barrier",
true)
2174 .
Case(
"bar.sync",
true)
2175 .
Case(
"barrier0.popc",
true)
2176 .
Case(
"barrier0.and",
true)
2177 .
Case(
"barrier0.or",
true)
2178 .
Case(
"clz.ll",
true)
2179 .
Case(
"popc.ll",
true)
2181 .
Case(
"swap.lo.hi.b64",
true)
2182 .
Case(
"tanh.approx.f32",
true)
2194 if (Name.starts_with(
"objectsize.")) {
2195 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2196 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2199 Intrinsic::objectsize, Tys);
2206 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2209 F->getParent(), Intrinsic::ptr_annotation,
2210 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2216 if (Name.consume_front(
"riscv.")) {
2219 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2220 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2221 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2222 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2225 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2238 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2239 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2248 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2249 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2250 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2251 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2256 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2265 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2267 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2276 if (Name ==
"stackprotectorcheck") {
2283 if (Name ==
"thread.pointer") {
2285 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2291 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2294 F->getParent(), Intrinsic::var_annotation,
2295 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2298 if (Name.consume_front(
"vector.splice")) {
2299 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2309 if (Name.consume_front(
"wasm.")) {
2312 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2313 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2314 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2319 F->getReturnType());
2323 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2325 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2327 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2346 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2356 std::string
Name =
F->getName().str();
2359 Name,
F->getParent());
2370 if (Result != std::nullopt) {
2387 bool CanUpgradeDebugIntrinsicsToRecords) {
2407 GV->
getName() ==
"llvm.global_dtors")) ||
2422 unsigned N =
Init->getNumOperands();
2423 std::vector<Constant *> NewCtors(
N);
2424 for (
unsigned i = 0; i !=
N; ++i) {
2427 Ctor->getAggregateElement(1),
2441 unsigned NumElts = ResultTy->getNumElements() * 8;
2445 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2455 for (
unsigned l = 0; l != NumElts; l += 16)
2456 for (
unsigned i = 0; i != 16; ++i) {
2457 unsigned Idx = NumElts + i - Shift;
2459 Idx -= NumElts - 16;
2460 Idxs[l + i] = Idx + l;
2463 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2467 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2475 unsigned NumElts = ResultTy->getNumElements() * 8;
2479 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2489 for (
unsigned l = 0; l != NumElts; l += 16)
2490 for (
unsigned i = 0; i != 16; ++i) {
2491 unsigned Idx = i + Shift;
2493 Idx += NumElts - 16;
2494 Idxs[l + i] = Idx + l;
2497 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2501 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2509 Mask = Builder.CreateBitCast(Mask, MaskTy);
2515 for (
unsigned i = 0; i != NumElts; ++i)
2517 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2528 if (
C->isAllOnesValue())
2533 return Builder.CreateSelect(Mask, Op0, Op1);
2540 if (
C->isAllOnesValue())
2544 Mask->getType()->getIntegerBitWidth());
2545 Mask = Builder.CreateBitCast(Mask, MaskTy);
2546 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2547 return Builder.CreateSelect(Mask, Op0, Op1);
2560 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2561 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2566 ShiftVal &= (NumElts - 1);
2575 if (ShiftVal > 16) {
2583 for (
unsigned l = 0; l < NumElts; l += 16) {
2584 for (
unsigned i = 0; i != 16; ++i) {
2585 unsigned Idx = ShiftVal + i;
2586 if (!IsVALIGN && Idx >= 16)
2587 Idx += NumElts - 16;
2588 Indices[l + i] = Idx + l;
2593 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2599 bool ZeroMask,
bool IndexForm) {
2602 unsigned EltWidth = Ty->getScalarSizeInBits();
2603 bool IsFloat = Ty->isFPOrFPVectorTy();
2605 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2606 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2607 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2608 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2609 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2610 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2611 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2612 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2613 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2614 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2615 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2616 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2617 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2618 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2619 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2620 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2621 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2622 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2623 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2624 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2625 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2626 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2627 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2628 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2629 else if (VecWidth == 128 && EltWidth == 16)
2630 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2631 else if (VecWidth == 256 && EltWidth == 16)
2632 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2633 else if (VecWidth == 512 && EltWidth == 16)
2634 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2635 else if (VecWidth == 128 && EltWidth == 8)
2636 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2637 else if (VecWidth == 256 && EltWidth == 8)
2638 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2639 else if (VecWidth == 512 && EltWidth == 8)
2640 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2651 Value *V = Builder.CreateIntrinsic(IID, Args);
2663 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2674 bool IsRotateRight) {
2684 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2685 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2688 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2689 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2734 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2739 bool IsShiftRight,
bool ZeroMask) {
2753 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2754 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2757 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2758 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2773 const Align Alignment =
2775 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2780 if (
C->isAllOnesValue())
2781 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2786 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2792 const Align Alignment =
2801 if (
C->isAllOnesValue())
2802 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2807 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2813 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2814 {Op0, Builder.getInt1(
false)});
2829 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2830 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2831 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2832 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2833 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2836 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2837 LHS = Builder.CreateAnd(
LHS, Mask);
2838 RHS = Builder.CreateAnd(
RHS, Mask);
2855 if (!
C || !
C->isAllOnesValue())
2856 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2861 for (
unsigned i = 0; i != NumElts; ++i)
2863 for (
unsigned i = NumElts; i != 8; ++i)
2864 Indices[i] = NumElts + i % NumElts;
2865 Vec = Builder.CreateShuffleVector(Vec,
2869 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2873 unsigned CC,
bool Signed) {
2881 }
else if (CC == 7) {
2917 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2918 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2920 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2921 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2930 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2936 Name = Name.substr(12);
2941 if (Name.starts_with(
"max.p")) {
2942 if (VecWidth == 128 && EltWidth == 32)
2943 IID = Intrinsic::x86_sse_max_ps;
2944 else if (VecWidth == 128 && EltWidth == 64)
2945 IID = Intrinsic::x86_sse2_max_pd;
2946 else if (VecWidth == 256 && EltWidth == 32)
2947 IID = Intrinsic::x86_avx_max_ps_256;
2948 else if (VecWidth == 256 && EltWidth == 64)
2949 IID = Intrinsic::x86_avx_max_pd_256;
2952 }
else if (Name.starts_with(
"min.p")) {
2953 if (VecWidth == 128 && EltWidth == 32)
2954 IID = Intrinsic::x86_sse_min_ps;
2955 else if (VecWidth == 128 && EltWidth == 64)
2956 IID = Intrinsic::x86_sse2_min_pd;
2957 else if (VecWidth == 256 && EltWidth == 32)
2958 IID = Intrinsic::x86_avx_min_ps_256;
2959 else if (VecWidth == 256 && EltWidth == 64)
2960 IID = Intrinsic::x86_avx_min_pd_256;
2963 }
else if (Name.starts_with(
"pshuf.b.")) {
2964 if (VecWidth == 128)
2965 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2966 else if (VecWidth == 256)
2967 IID = Intrinsic::x86_avx2_pshuf_b;
2968 else if (VecWidth == 512)
2969 IID = Intrinsic::x86_avx512_pshuf_b_512;
2972 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2973 if (VecWidth == 128)
2974 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2975 else if (VecWidth == 256)
2976 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2977 else if (VecWidth == 512)
2978 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2981 }
else if (Name.starts_with(
"pmulh.w.")) {
2982 if (VecWidth == 128)
2983 IID = Intrinsic::x86_sse2_pmulh_w;
2984 else if (VecWidth == 256)
2985 IID = Intrinsic::x86_avx2_pmulh_w;
2986 else if (VecWidth == 512)
2987 IID = Intrinsic::x86_avx512_pmulh_w_512;
2990 }
else if (Name.starts_with(
"pmulhu.w.")) {
2991 if (VecWidth == 128)
2992 IID = Intrinsic::x86_sse2_pmulhu_w;
2993 else if (VecWidth == 256)
2994 IID = Intrinsic::x86_avx2_pmulhu_w;
2995 else if (VecWidth == 512)
2996 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2999 }
else if (Name.starts_with(
"pmaddw.d.")) {
3000 if (VecWidth == 128)
3001 IID = Intrinsic::x86_sse2_pmadd_wd;
3002 else if (VecWidth == 256)
3003 IID = Intrinsic::x86_avx2_pmadd_wd;
3004 else if (VecWidth == 512)
3005 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3008 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3009 if (VecWidth == 128)
3010 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3011 else if (VecWidth == 256)
3012 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3013 else if (VecWidth == 512)
3014 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3017 }
else if (Name.starts_with(
"packsswb.")) {
3018 if (VecWidth == 128)
3019 IID = Intrinsic::x86_sse2_packsswb_128;
3020 else if (VecWidth == 256)
3021 IID = Intrinsic::x86_avx2_packsswb;
3022 else if (VecWidth == 512)
3023 IID = Intrinsic::x86_avx512_packsswb_512;
3026 }
else if (Name.starts_with(
"packssdw.")) {
3027 if (VecWidth == 128)
3028 IID = Intrinsic::x86_sse2_packssdw_128;
3029 else if (VecWidth == 256)
3030 IID = Intrinsic::x86_avx2_packssdw;
3031 else if (VecWidth == 512)
3032 IID = Intrinsic::x86_avx512_packssdw_512;
3035 }
else if (Name.starts_with(
"packuswb.")) {
3036 if (VecWidth == 128)
3037 IID = Intrinsic::x86_sse2_packuswb_128;
3038 else if (VecWidth == 256)
3039 IID = Intrinsic::x86_avx2_packuswb;
3040 else if (VecWidth == 512)
3041 IID = Intrinsic::x86_avx512_packuswb_512;
3044 }
else if (Name.starts_with(
"packusdw.")) {
3045 if (VecWidth == 128)
3046 IID = Intrinsic::x86_sse41_packusdw;
3047 else if (VecWidth == 256)
3048 IID = Intrinsic::x86_avx2_packusdw;
3049 else if (VecWidth == 512)
3050 IID = Intrinsic::x86_avx512_packusdw_512;
3053 }
else if (Name.starts_with(
"vpermilvar.")) {
3054 if (VecWidth == 128 && EltWidth == 32)
3055 IID = Intrinsic::x86_avx_vpermilvar_ps;
3056 else if (VecWidth == 128 && EltWidth == 64)
3057 IID = Intrinsic::x86_avx_vpermilvar_pd;
3058 else if (VecWidth == 256 && EltWidth == 32)
3059 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3060 else if (VecWidth == 256 && EltWidth == 64)
3061 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3062 else if (VecWidth == 512 && EltWidth == 32)
3063 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3064 else if (VecWidth == 512 && EltWidth == 64)
3065 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3068 }
else if (Name ==
"cvtpd2dq.256") {
3069 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3070 }
else if (Name ==
"cvtpd2ps.256") {
3071 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3072 }
else if (Name ==
"cvttpd2dq.256") {
3073 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3074 }
else if (Name ==
"cvttps2dq.128") {
3075 IID = Intrinsic::x86_sse2_cvttps2dq;
3076 }
else if (Name ==
"cvttps2dq.256") {
3077 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3078 }
else if (Name.starts_with(
"permvar.")) {
3080 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3081 IID = Intrinsic::x86_avx2_permps;
3082 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3083 IID = Intrinsic::x86_avx2_permd;
3084 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3085 IID = Intrinsic::x86_avx512_permvar_df_256;
3086 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3087 IID = Intrinsic::x86_avx512_permvar_di_256;
3088 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3089 IID = Intrinsic::x86_avx512_permvar_sf_512;
3090 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3091 IID = Intrinsic::x86_avx512_permvar_si_512;
3092 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3093 IID = Intrinsic::x86_avx512_permvar_df_512;
3094 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3095 IID = Intrinsic::x86_avx512_permvar_di_512;
3096 else if (VecWidth == 128 && EltWidth == 16)
3097 IID = Intrinsic::x86_avx512_permvar_hi_128;
3098 else if (VecWidth == 256 && EltWidth == 16)
3099 IID = Intrinsic::x86_avx512_permvar_hi_256;
3100 else if (VecWidth == 512 && EltWidth == 16)
3101 IID = Intrinsic::x86_avx512_permvar_hi_512;
3102 else if (VecWidth == 128 && EltWidth == 8)
3103 IID = Intrinsic::x86_avx512_permvar_qi_128;
3104 else if (VecWidth == 256 && EltWidth == 8)
3105 IID = Intrinsic::x86_avx512_permvar_qi_256;
3106 else if (VecWidth == 512 && EltWidth == 8)
3107 IID = Intrinsic::x86_avx512_permvar_qi_512;
3110 }
else if (Name.starts_with(
"dbpsadbw.")) {
3111 if (VecWidth == 128)
3112 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3113 else if (VecWidth == 256)
3114 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3115 else if (VecWidth == 512)
3116 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3119 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3120 if (VecWidth == 128)
3121 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3122 else if (VecWidth == 256)
3123 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3124 else if (VecWidth == 512)
3125 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3128 }
else if (Name.starts_with(
"conflict.")) {
3129 if (Name[9] ==
'd' && VecWidth == 128)
3130 IID = Intrinsic::x86_avx512_conflict_d_128;
3131 else if (Name[9] ==
'd' && VecWidth == 256)
3132 IID = Intrinsic::x86_avx512_conflict_d_256;
3133 else if (Name[9] ==
'd' && VecWidth == 512)
3134 IID = Intrinsic::x86_avx512_conflict_d_512;
3135 else if (Name[9] ==
'q' && VecWidth == 128)
3136 IID = Intrinsic::x86_avx512_conflict_q_128;
3137 else if (Name[9] ==
'q' && VecWidth == 256)
3138 IID = Intrinsic::x86_avx512_conflict_q_256;
3139 else if (Name[9] ==
'q' && VecWidth == 512)
3140 IID = Intrinsic::x86_avx512_conflict_q_512;
3143 }
else if (Name.starts_with(
"pavg.")) {
3144 if (Name[5] ==
'b' && VecWidth == 128)
3145 IID = Intrinsic::x86_sse2_pavg_b;
3146 else if (Name[5] ==
'b' && VecWidth == 256)
3147 IID = Intrinsic::x86_avx2_pavg_b;
3148 else if (Name[5] ==
'b' && VecWidth == 512)
3149 IID = Intrinsic::x86_avx512_pavg_b_512;
3150 else if (Name[5] ==
'w' && VecWidth == 128)
3151 IID = Intrinsic::x86_sse2_pavg_w;
3152 else if (Name[5] ==
'w' && VecWidth == 256)
3153 IID = Intrinsic::x86_avx2_pavg_w;
3154 else if (Name[5] ==
'w' && VecWidth == 512)
3155 IID = Intrinsic::x86_avx512_pavg_w_512;
3164 Rep = Builder.CreateIntrinsic(IID, Args);
3175 if (AsmStr->find(
"mov\tfp") == 0 &&
3176 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3177 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3178 AsmStr->replace(Pos, 1,
";");
3184 Value *Rep =
nullptr;
3186 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3188 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3189 {Arg, Builder.getTrue()},
3191 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3192 Type *Ty = (Name ==
"abs.bf16")
3196 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3197 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3198 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3199 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3200 : Intrinsic::nvvm_fabs;
3201 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3202 }
else if (Name.consume_front(
"add.")) {
3205 assert(
FAdd &&
"unsupported nvvm.add.* intrinsic");
3208 Rep = Builder.CreateIntrinsic(
3210 {A, CI->getArgOperand(1),
3211 Builder.getInt32(static_cast<int>(RoundingMode))});
3212 }
else if (Name.consume_front(
"ex2.approx.")) {
3214 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3215 : Intrinsic::nvvm_ex2_approx;
3216 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3217 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3218 Name.starts_with(
"atomic.load.add.f64.p")) {
3221 Rep = Builder.CreateAtomicRMW(
3227 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3228 Name.starts_with(
"atomic.load.dec.32.p")) {
3233 Rep = Builder.CreateAtomicRMW(
3237 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3243 Op.contains(
".cta.") ?
"block" :
"");
3244 if (
Op.starts_with(
"cas.")) {
3246 Value *Pair = Builder.CreateAtomicCmpXchg(
3249 Rep = Builder.CreateExtractValue(Pair, 0);
3267 "unexpected nvvm scoped atomic intrinsic");
3268 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3271 }
else if (Name ==
"clz.ll") {
3274 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3275 {Arg, Builder.getFalse()},
3277 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3278 }
else if (Name ==
"popc.ll") {
3282 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3283 Arg,
nullptr,
"ctpop");
3284 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3285 }
else if (Name ==
"h2f") {
3287 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3288 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3289 }
else if (Name.consume_front(
"bitcast.") &&
3290 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3293 }
else if (Name ==
"rotate.b32") {
3296 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3297 {Arg, Arg, ShiftAmt});
3298 }
else if (Name ==
"rotate.b64") {
3302 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3303 {Arg, Arg, ZExtShiftAmt});
3304 }
else if (Name ==
"rotate.right.b64") {
3308 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3309 {Arg, Arg, ZExtShiftAmt});
3310 }
else if (Name ==
"swap.lo.hi.b64") {
3313 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3314 {Arg, Arg, Builder.getInt64(32)});
3315 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3318 Name.starts_with(
".to.gen"))) {
3320 }
else if (Name.consume_front(
"ldg.global")) {
3324 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3327 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3329 }
else if (Name ==
"tanh.approx.f32") {
3333 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3335 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3337 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3338 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3340 }
else if (Name ==
"barrier") {
3341 Rep = Builder.CreateIntrinsic(
3342 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3344 }
else if (Name ==
"barrier.sync") {
3345 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3347 }
else if (Name ==
"barrier.sync.cnt") {
3348 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3350 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3351 Name ==
"barrier0.or") {
3353 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3357 .
Case(
"barrier0.popc",
3358 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3359 .
Case(
"barrier0.and",
3360 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3361 .
Case(
"barrier0.or",
3362 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3363 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3364 Rep = Builder.CreateZExt(Bar, CI->
getType());
3378 ? Builder.CreateBitCast(Arg, NewType)
3381 Rep = Builder.CreateCall(NewFn, Args);
3382 if (
F->getReturnType()->isIntegerTy())
3383 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3393 Value *Rep =
nullptr;
3395 if (Name.starts_with(
"sse4a.movnt.")) {
3407 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3410 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3411 }
else if (Name.starts_with(
"avx.movnt.") ||
3412 Name.starts_with(
"avx512.storent.")) {
3424 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3425 }
else if (Name ==
"sse2.storel.dq") {
3430 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3431 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3432 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3433 }
else if (Name.starts_with(
"sse.storeu.") ||
3434 Name.starts_with(
"sse2.storeu.") ||
3435 Name.starts_with(
"avx.storeu.")) {
3438 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3439 }
else if (Name ==
"avx512.mask.store.ss") {
3443 }
else if (Name.starts_with(
"avx512.mask.store")) {
3445 bool Aligned = Name[17] !=
'u';
3448 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3451 bool CmpEq = Name[9] ==
'e';
3454 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3455 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3462 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3463 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3465 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3466 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3467 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3468 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3469 Name.starts_with(
"sse2.sqrt.p") ||
3470 Name.starts_with(
"sse.sqrt.p")) {
3471 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3472 {CI->getArgOperand(0)});
3473 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3477 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3478 : Intrinsic::x86_avx512_sqrt_pd_512;
3481 Rep = Builder.CreateIntrinsic(IID, Args);
3483 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3484 {CI->getArgOperand(0)});
3488 }
else if (Name.starts_with(
"avx512.ptestm") ||
3489 Name.starts_with(
"avx512.ptestnm")) {
3493 Rep = Builder.CreateAnd(Op0, Op1);
3499 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3501 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3504 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3507 }
else if (Name.starts_with(
"avx512.kunpck")) {
3512 for (
unsigned i = 0; i != NumElts; ++i)
3521 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3522 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3523 }
else if (Name ==
"avx512.kand.w") {
3526 Rep = Builder.CreateAnd(
LHS,
RHS);
3527 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3528 }
else if (Name ==
"avx512.kandn.w") {
3531 LHS = Builder.CreateNot(
LHS);
3532 Rep = Builder.CreateAnd(
LHS,
RHS);
3533 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3534 }
else if (Name ==
"avx512.kor.w") {
3537 Rep = Builder.CreateOr(
LHS,
RHS);
3538 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3539 }
else if (Name ==
"avx512.kxor.w") {
3542 Rep = Builder.CreateXor(
LHS,
RHS);
3543 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3544 }
else if (Name ==
"avx512.kxnor.w") {
3547 LHS = Builder.CreateNot(
LHS);
3548 Rep = Builder.CreateXor(
LHS,
RHS);
3549 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3550 }
else if (Name ==
"avx512.knot.w") {
3552 Rep = Builder.CreateNot(Rep);
3553 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3554 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3557 Rep = Builder.CreateOr(
LHS,
RHS);
3558 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3560 if (Name[14] ==
'c')
3564 Rep = Builder.CreateICmpEQ(Rep,
C);
3565 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3566 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3567 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3568 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3569 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3572 ConstantInt::get(I32Ty, 0));
3574 ConstantInt::get(I32Ty, 0));
3576 if (Name.contains(
".add."))
3577 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3578 else if (Name.contains(
".sub."))
3579 EltOp = Builder.CreateFSub(Elt0, Elt1);
3580 else if (Name.contains(
".mul."))
3581 EltOp = Builder.CreateFMul(Elt0, Elt1);
3583 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3584 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3585 ConstantInt::get(I32Ty, 0));
3586 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3588 bool CmpEq = Name[16] ==
'e';
3590 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3592 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3599 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3602 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3605 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3612 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3614 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3615 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3617 if (VecWidth == 128 && EltWidth == 32)
3618 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3619 else if (VecWidth == 256 && EltWidth == 32)
3620 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3621 else if (VecWidth == 512 && EltWidth == 32)
3622 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3623 else if (VecWidth == 128 && EltWidth == 64)
3624 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3625 else if (VecWidth == 256 && EltWidth == 64)
3626 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3627 else if (VecWidth == 512 && EltWidth == 64)
3628 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3635 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3638 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3639 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3641 if (VecWidth == 128 && EltWidth == 32)
3642 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3643 else if (VecWidth == 256 && EltWidth == 32)
3644 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3645 else if (VecWidth == 512 && EltWidth == 32)
3646 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3647 else if (VecWidth == 128 && EltWidth == 64)
3648 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3649 else if (VecWidth == 256 && EltWidth == 64)
3650 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3651 else if (VecWidth == 512 && EltWidth == 64)
3652 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3657 if (VecWidth == 512)
3659 Args.push_back(Mask);
3661 Rep = Builder.CreateIntrinsic(IID, Args);
3662 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3666 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3669 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3670 Name.starts_with(
"avx512.cvtw2mask.") ||
3671 Name.starts_with(
"avx512.cvtd2mask.") ||
3672 Name.starts_with(
"avx512.cvtq2mask.")) {
3677 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3678 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3679 Name.starts_with(
"avx512.mask.pabs")) {
3681 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3682 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3683 Name.starts_with(
"avx512.mask.pmaxs")) {
3685 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3686 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3687 Name.starts_with(
"avx512.mask.pmaxu")) {
3689 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3690 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3691 Name.starts_with(
"avx512.mask.pmins")) {
3693 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3694 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3695 Name.starts_with(
"avx512.mask.pminu")) {
3697 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3698 Name ==
"avx512.pmulu.dq.512" ||
3699 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3701 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3702 Name ==
"avx512.pmul.dq.512" ||
3703 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3705 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3706 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3711 }
else if (Name ==
"avx512.cvtusi2sd") {
3716 }
else if (Name ==
"sse2.cvtss2sd") {
3718 Rep = Builder.CreateFPExt(
3721 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3722 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3723 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3724 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3725 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3726 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3727 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3728 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3729 Name ==
"avx512.mask.cvtqq2ps.256" ||
3730 Name ==
"avx512.mask.cvtqq2ps.512" ||
3731 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3732 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3733 Name ==
"avx.cvt.ps2.pd.256" ||
3734 Name ==
"avx512.mask.cvtps2pd.128" ||
3735 Name ==
"avx512.mask.cvtps2pd.256") {
3740 unsigned NumDstElts = DstTy->getNumElements();
3741 if (NumDstElts < SrcTy->getNumElements()) {
3742 assert(NumDstElts == 2 &&
"Unexpected vector size");
3743 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3746 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3747 bool IsUnsigned = Name.contains(
"cvtu");
3749 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3753 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3754 : Intrinsic::x86_avx512_sitofp_round;
3755 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3758 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3759 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3765 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3766 Name.starts_with(
"vcvtph2ps.")) {
3770 unsigned NumDstElts = DstTy->getNumElements();
3771 if (NumDstElts != SrcTy->getNumElements()) {
3772 assert(NumDstElts == 4 &&
"Unexpected vector size");
3773 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3775 Rep = Builder.CreateBitCast(
3777 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3781 }
else if (Name.starts_with(
"avx512.mask.load")) {
3783 bool Aligned = Name[16] !=
'u';
3786 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3790 ResultTy->getNumElements());
3791 Rep = Builder.CreateIntrinsic(
3792 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3794 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3800 Rep = Builder.CreateIntrinsic(
3801 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3803 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3804 Name.starts_with(
"avx512.mask.expand.")) {
3808 ResultTy->getNumElements());
3810 bool IsCompress = Name[12] ==
'c';
3811 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3812 : Intrinsic::x86_avx512_mask_expand;
3813 Rep = Builder.CreateIntrinsic(
3815 }
else if (Name.starts_with(
"xop.vpcom")) {
3817 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3818 Name.ends_with(
"uq"))
3820 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3821 Name.ends_with(
"d") || Name.ends_with(
"q"))
3830 Name = Name.substr(9);
3831 if (Name.starts_with(
"lt"))
3833 else if (Name.starts_with(
"le"))
3835 else if (Name.starts_with(
"gt"))
3837 else if (Name.starts_with(
"ge"))
3839 else if (Name.starts_with(
"eq"))
3841 else if (Name.starts_with(
"ne"))
3843 else if (Name.starts_with(
"false"))
3845 else if (Name.starts_with(
"true"))
3852 }
else if (Name.starts_with(
"xop.vpcmov")) {
3854 Value *NotSel = Builder.CreateNot(Sel);
3857 Rep = Builder.CreateOr(Sel0, Sel1);
3858 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3859 Name.starts_with(
"avx512.mask.prol")) {
3861 }
else if (Name.starts_with(
"avx512.pror") ||
3862 Name.starts_with(
"avx512.mask.pror")) {
3864 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3865 Name.starts_with(
"avx512.mask.vpshld") ||
3866 Name.starts_with(
"avx512.maskz.vpshld")) {
3867 bool ZeroMask = Name[11] ==
'z';
3869 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3870 Name.starts_with(
"avx512.mask.vpshrd") ||
3871 Name.starts_with(
"avx512.maskz.vpshrd")) {
3872 bool ZeroMask = Name[11] ==
'z';
3874 }
else if (Name ==
"sse42.crc32.64.8") {
3877 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3879 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3880 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3881 Name.starts_with(
"avx512.vbroadcast.s")) {
3884 Type *EltTy = VecTy->getElementType();
3885 unsigned EltNum = VecTy->getNumElements();
3889 for (
unsigned I = 0;
I < EltNum; ++
I)
3890 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3891 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3892 Name.starts_with(
"sse41.pmovzx") ||
3893 Name.starts_with(
"avx2.pmovsx") ||
3894 Name.starts_with(
"avx2.pmovzx") ||
3895 Name.starts_with(
"avx512.mask.pmovsx") ||
3896 Name.starts_with(
"avx512.mask.pmovzx")) {
3898 unsigned NumDstElts = DstTy->getNumElements();
3902 for (
unsigned i = 0; i != NumDstElts; ++i)
3907 bool DoSext = Name.contains(
"pmovsx");
3909 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3914 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3915 Name ==
"avx512.mask.pmov.qd.512" ||
3916 Name ==
"avx512.mask.pmov.wb.256" ||
3917 Name ==
"avx512.mask.pmov.wb.512") {
3922 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3923 Name ==
"avx2.vbroadcasti128") {
3929 if (NumSrcElts == 2)
3932 Rep = Builder.CreateShuffleVector(
Load,
3934 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3935 Name.starts_with(
"avx512.mask.shuf.f")) {
3940 unsigned ControlBitsMask = NumLanes - 1;
3941 unsigned NumControlBits = NumLanes / 2;
3944 for (
unsigned l = 0; l != NumLanes; ++l) {
3945 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
3947 if (l >= NumLanes / 2)
3948 LaneMask += NumLanes;
3949 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3950 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3956 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3957 Name.starts_with(
"avx512.mask.broadcasti")) {
3960 unsigned NumDstElts =
3964 for (
unsigned i = 0; i != NumDstElts; ++i)
3965 ShuffleMask[i] = i % NumSrcElts;
3971 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3972 Name.starts_with(
"avx2.vbroadcast") ||
3973 Name.starts_with(
"avx512.pbroadcast") ||
3974 Name.starts_with(
"avx512.mask.broadcast.s")) {
3981 Rep = Builder.CreateShuffleVector(
Op, M);
3986 }
else if (Name.starts_with(
"sse2.padds.") ||
3987 Name.starts_with(
"avx2.padds.") ||
3988 Name.starts_with(
"avx512.padds.") ||
3989 Name.starts_with(
"avx512.mask.padds.")) {
3991 }
else if (Name.starts_with(
"sse2.psubs.") ||
3992 Name.starts_with(
"avx2.psubs.") ||
3993 Name.starts_with(
"avx512.psubs.") ||
3994 Name.starts_with(
"avx512.mask.psubs.")) {
3996 }
else if (Name.starts_with(
"sse2.paddus.") ||
3997 Name.starts_with(
"avx2.paddus.") ||
3998 Name.starts_with(
"avx512.mask.paddus.")) {
4000 }
else if (Name.starts_with(
"sse2.psubus.") ||
4001 Name.starts_with(
"avx2.psubus.") ||
4002 Name.starts_with(
"avx512.mask.psubus.")) {
4004 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4009 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4013 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4018 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4023 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4024 Name ==
"avx512.psll.dq.512") {
4028 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4029 Name ==
"avx512.psrl.dq.512") {
4033 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4034 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4035 Name.starts_with(
"avx2.pblendd.")) {
4040 unsigned NumElts = VecTy->getNumElements();
4043 for (
unsigned i = 0; i != NumElts; ++i)
4044 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4046 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4047 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4048 Name ==
"avx2.vinserti128" ||
4049 Name.starts_with(
"avx512.mask.insert")) {
4053 unsigned DstNumElts =
4055 unsigned SrcNumElts =
4057 unsigned Scale = DstNumElts / SrcNumElts;
4064 for (
unsigned i = 0; i != SrcNumElts; ++i)
4066 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4067 Idxs[i] = SrcNumElts;
4068 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4082 for (
unsigned i = 0; i != DstNumElts; ++i)
4085 for (
unsigned i = 0; i != SrcNumElts; ++i)
4086 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4087 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4093 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4094 Name ==
"avx2.vextracti128" ||
4095 Name.starts_with(
"avx512.mask.vextract")) {
4098 unsigned DstNumElts =
4100 unsigned SrcNumElts =
4102 unsigned Scale = SrcNumElts / DstNumElts;
4109 for (
unsigned i = 0; i != DstNumElts; ++i) {
4110 Idxs[i] = i + (
Imm * DstNumElts);
4112 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4118 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4119 Name.starts_with(
"avx512.mask.perm.di.")) {
4123 unsigned NumElts = VecTy->getNumElements();
4126 for (
unsigned i = 0; i != NumElts; ++i)
4127 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4129 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4134 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4146 unsigned HalfSize = NumElts / 2;
4158 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4159 for (
unsigned i = 0; i < HalfSize; ++i)
4160 ShuffleMask[i] = StartIndex + i;
4163 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4164 for (
unsigned i = 0; i < HalfSize; ++i)
4165 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4167 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4169 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4170 Name.starts_with(
"avx512.mask.vpermil.p") ||
4171 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4175 unsigned NumElts = VecTy->getNumElements();
4177 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4178 unsigned IdxMask = ((1 << IdxSize) - 1);
4184 for (
unsigned i = 0; i != NumElts; ++i)
4185 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4187 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4192 }
else if (Name ==
"sse2.pshufl.w" ||
4193 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4198 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4202 for (
unsigned l = 0; l != NumElts; l += 8) {
4203 for (
unsigned i = 0; i != 4; ++i)
4204 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4205 for (
unsigned i = 4; i != 8; ++i)
4206 Idxs[i + l] = i + l;
4209 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4214 }
else if (Name ==
"sse2.pshufh.w" ||
4215 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4220 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4224 for (
unsigned l = 0; l != NumElts; l += 8) {
4225 for (
unsigned i = 0; i != 4; ++i)
4226 Idxs[i + l] = i + l;
4227 for (
unsigned i = 0; i != 4; ++i)
4228 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4231 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4236 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4243 unsigned HalfLaneElts = NumLaneElts / 2;
4246 for (
unsigned i = 0; i != NumElts; ++i) {
4248 Idxs[i] = i - (i % NumLaneElts);
4250 if ((i % NumLaneElts) >= HalfLaneElts)
4254 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4257 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4261 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4262 Name.starts_with(
"avx512.mask.movshdup") ||
4263 Name.starts_with(
"avx512.mask.movsldup")) {
4269 if (Name.starts_with(
"avx512.mask.movshdup."))
4273 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4274 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4275 Idxs[i + l + 0] = i + l +
Offset;
4276 Idxs[i + l + 1] = i + l +
Offset;
4279 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4283 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4284 Name.starts_with(
"avx512.mask.unpckl.")) {
4291 for (
int l = 0; l != NumElts; l += NumLaneElts)
4292 for (
int i = 0; i != NumLaneElts; ++i)
4293 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4295 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4299 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4300 Name.starts_with(
"avx512.mask.unpckh.")) {
4307 for (
int l = 0; l != NumElts; l += NumLaneElts)
4308 for (
int i = 0; i != NumLaneElts; ++i)
4309 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4311 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4315 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4316 Name.starts_with(
"avx512.mask.pand.")) {
4319 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4321 Rep = Builder.CreateBitCast(Rep, FTy);
4324 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4325 Name.starts_with(
"avx512.mask.pandn.")) {
4328 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4329 Rep = Builder.CreateAnd(Rep,
4331 Rep = Builder.CreateBitCast(Rep, FTy);
4334 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4335 Name.starts_with(
"avx512.mask.por.")) {
4338 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4340 Rep = Builder.CreateBitCast(Rep, FTy);
4343 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4344 Name.starts_with(
"avx512.mask.pxor.")) {
4347 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4349 Rep = Builder.CreateBitCast(Rep, FTy);
4352 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4356 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4360 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4364 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4365 if (Name.ends_with(
".512")) {
4367 if (Name[17] ==
's')
4368 IID = Intrinsic::x86_avx512_add_ps_512;
4370 IID = Intrinsic::x86_avx512_add_pd_512;
4372 Rep = Builder.CreateIntrinsic(
4380 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4381 if (Name.ends_with(
".512")) {
4383 if (Name[17] ==
's')
4384 IID = Intrinsic::x86_avx512_div_ps_512;
4386 IID = Intrinsic::x86_avx512_div_pd_512;
4388 Rep = Builder.CreateIntrinsic(
4396 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4397 if (Name.ends_with(
".512")) {
4399 if (Name[17] ==
's')
4400 IID = Intrinsic::x86_avx512_mul_ps_512;
4402 IID = Intrinsic::x86_avx512_mul_pd_512;
4404 Rep = Builder.CreateIntrinsic(
4412 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4413 if (Name.ends_with(
".512")) {
4415 if (Name[17] ==
's')
4416 IID = Intrinsic::x86_avx512_sub_ps_512;
4418 IID = Intrinsic::x86_avx512_sub_pd_512;
4420 Rep = Builder.CreateIntrinsic(
4428 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4429 Name.starts_with(
"avx512.mask.min.p")) &&
4430 Name.drop_front(18) ==
".512") {
4431 bool IsDouble = Name[17] ==
'd';
4432 bool IsMin = Name[13] ==
'i';
4434 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4435 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4438 Rep = Builder.CreateIntrinsic(
4443 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4445 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4446 {CI->getArgOperand(0), Builder.getInt1(false)});
4449 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4450 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4451 bool IsVariable = Name[16] ==
'v';
4452 char Size = Name[16] ==
'.' ? Name[17]
4453 : Name[17] ==
'.' ? Name[18]
4454 : Name[18] ==
'.' ? Name[19]
4458 if (IsVariable && Name[17] !=
'.') {
4459 if (
Size ==
'd' && Name[17] ==
'2')
4460 IID = Intrinsic::x86_avx2_psllv_q;
4461 else if (
Size ==
'd' && Name[17] ==
'4')
4462 IID = Intrinsic::x86_avx2_psllv_q_256;
4463 else if (
Size ==
's' && Name[17] ==
'4')
4464 IID = Intrinsic::x86_avx2_psllv_d;
4465 else if (
Size ==
's' && Name[17] ==
'8')
4466 IID = Intrinsic::x86_avx2_psllv_d_256;
4467 else if (
Size ==
'h' && Name[17] ==
'8')
4468 IID = Intrinsic::x86_avx512_psllv_w_128;
4469 else if (
Size ==
'h' && Name[17] ==
'1')
4470 IID = Intrinsic::x86_avx512_psllv_w_256;
4471 else if (Name[17] ==
'3' && Name[18] ==
'2')
4472 IID = Intrinsic::x86_avx512_psllv_w_512;
4475 }
else if (Name.ends_with(
".128")) {
4477 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4478 : Intrinsic::x86_sse2_psll_d;
4479 else if (
Size ==
'q')
4480 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4481 : Intrinsic::x86_sse2_psll_q;
4482 else if (
Size ==
'w')
4483 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4484 : Intrinsic::x86_sse2_psll_w;
4487 }
else if (Name.ends_with(
".256")) {
4489 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4490 : Intrinsic::x86_avx2_psll_d;
4491 else if (
Size ==
'q')
4492 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4493 : Intrinsic::x86_avx2_psll_q;
4494 else if (
Size ==
'w')
4495 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4496 : Intrinsic::x86_avx2_psll_w;
4501 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4502 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4503 : Intrinsic::x86_avx512_psll_d_512;
4504 else if (
Size ==
'q')
4505 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4506 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4507 : Intrinsic::x86_avx512_psll_q_512;
4508 else if (
Size ==
'w')
4509 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4510 : Intrinsic::x86_avx512_psll_w_512;
4516 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4517 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4518 bool IsVariable = Name[16] ==
'v';
4519 char Size = Name[16] ==
'.' ? Name[17]
4520 : Name[17] ==
'.' ? Name[18]
4521 : Name[18] ==
'.' ? Name[19]
4525 if (IsVariable && Name[17] !=
'.') {
4526 if (
Size ==
'd' && Name[17] ==
'2')
4527 IID = Intrinsic::x86_avx2_psrlv_q;
4528 else if (
Size ==
'd' && Name[17] ==
'4')
4529 IID = Intrinsic::x86_avx2_psrlv_q_256;
4530 else if (
Size ==
's' && Name[17] ==
'4')
4531 IID = Intrinsic::x86_avx2_psrlv_d;
4532 else if (
Size ==
's' && Name[17] ==
'8')
4533 IID = Intrinsic::x86_avx2_psrlv_d_256;
4534 else if (
Size ==
'h' && Name[17] ==
'8')
4535 IID = Intrinsic::x86_avx512_psrlv_w_128;
4536 else if (
Size ==
'h' && Name[17] ==
'1')
4537 IID = Intrinsic::x86_avx512_psrlv_w_256;
4538 else if (Name[17] ==
'3' && Name[18] ==
'2')
4539 IID = Intrinsic::x86_avx512_psrlv_w_512;
4542 }
else if (Name.ends_with(
".128")) {
4544 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4545 : Intrinsic::x86_sse2_psrl_d;
4546 else if (
Size ==
'q')
4547 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4548 : Intrinsic::x86_sse2_psrl_q;
4549 else if (
Size ==
'w')
4550 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4551 : Intrinsic::x86_sse2_psrl_w;
4554 }
else if (Name.ends_with(
".256")) {
4556 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4557 : Intrinsic::x86_avx2_psrl_d;
4558 else if (
Size ==
'q')
4559 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4560 : Intrinsic::x86_avx2_psrl_q;
4561 else if (
Size ==
'w')
4562 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4563 : Intrinsic::x86_avx2_psrl_w;
4568 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4569 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4570 : Intrinsic::x86_avx512_psrl_d_512;
4571 else if (
Size ==
'q')
4572 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4573 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4574 : Intrinsic::x86_avx512_psrl_q_512;
4575 else if (
Size ==
'w')
4576 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4577 : Intrinsic::x86_avx512_psrl_w_512;
4583 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4584 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4585 bool IsVariable = Name[16] ==
'v';
4586 char Size = Name[16] ==
'.' ? Name[17]
4587 : Name[17] ==
'.' ? Name[18]
4588 : Name[18] ==
'.' ? Name[19]
4592 if (IsVariable && Name[17] !=
'.') {
4593 if (
Size ==
's' && Name[17] ==
'4')
4594 IID = Intrinsic::x86_avx2_psrav_d;
4595 else if (
Size ==
's' && Name[17] ==
'8')
4596 IID = Intrinsic::x86_avx2_psrav_d_256;
4597 else if (
Size ==
'h' && Name[17] ==
'8')
4598 IID = Intrinsic::x86_avx512_psrav_w_128;
4599 else if (
Size ==
'h' && Name[17] ==
'1')
4600 IID = Intrinsic::x86_avx512_psrav_w_256;
4601 else if (Name[17] ==
'3' && Name[18] ==
'2')
4602 IID = Intrinsic::x86_avx512_psrav_w_512;
4605 }
else if (Name.ends_with(
".128")) {
4607 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4608 : Intrinsic::x86_sse2_psra_d;
4609 else if (
Size ==
'q')
4610 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4611 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4612 : Intrinsic::x86_avx512_psra_q_128;
4613 else if (
Size ==
'w')
4614 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4615 : Intrinsic::x86_sse2_psra_w;
4618 }
else if (Name.ends_with(
".256")) {
4620 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4621 : Intrinsic::x86_avx2_psra_d;
4622 else if (
Size ==
'q')
4623 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4624 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4625 : Intrinsic::x86_avx512_psra_q_256;
4626 else if (
Size ==
'w')
4627 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4628 : Intrinsic::x86_avx2_psra_w;
4633 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4634 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4635 : Intrinsic::x86_avx512_psra_d_512;
4636 else if (
Size ==
'q')
4637 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4638 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4639 : Intrinsic::x86_avx512_psra_q_512;
4640 else if (
Size ==
'w')
4641 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4642 : Intrinsic::x86_avx512_psra_w_512;
4648 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4650 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4652 }
else if (Name.ends_with(
".movntdqa")) {
4656 LoadInst *LI = Builder.CreateAlignedLoad(
4661 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4662 Name.starts_with(
"fma.vfmsub.") ||
4663 Name.starts_with(
"fma.vfnmadd.") ||
4664 Name.starts_with(
"fma.vfnmsub.")) {
4665 bool NegMul = Name[6] ==
'n';
4666 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4667 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4678 if (NegMul && !IsScalar)
4679 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4680 if (NegMul && IsScalar)
4681 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4683 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4685 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4689 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4697 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4701 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4702 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4703 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4704 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4705 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4706 bool IsMask3 = Name[11] ==
'3';
4707 bool IsMaskZ = Name[11] ==
'z';
4709 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4710 bool NegMul = Name[2] ==
'n';
4711 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4717 if (NegMul && (IsMask3 || IsMaskZ))
4718 A = Builder.CreateFNeg(
A);
4719 if (NegMul && !(IsMask3 || IsMaskZ))
4720 B = Builder.CreateFNeg(
B);
4722 C = Builder.CreateFNeg(
C);
4724 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4725 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4726 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4733 if (Name.back() ==
'd')
4734 IID = Intrinsic::x86_avx512_vfmadd_f64;
4736 IID = Intrinsic::x86_avx512_vfmadd_f32;
4737 Rep = Builder.CreateIntrinsic(IID,
Ops);
4739 Rep = Builder.CreateFMA(
A,
B,
C);
4748 if (NegAcc && IsMask3)
4753 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4755 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4756 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4757 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4758 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4759 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4760 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4761 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4762 bool IsMask3 = Name[11] ==
'3';
4763 bool IsMaskZ = Name[11] ==
'z';
4765 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4766 bool NegMul = Name[2] ==
'n';
4767 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4773 if (NegMul && (IsMask3 || IsMaskZ))
4774 A = Builder.CreateFNeg(
A);
4775 if (NegMul && !(IsMask3 || IsMaskZ))
4776 B = Builder.CreateFNeg(
B);
4778 C = Builder.CreateFNeg(
C);
4785 if (Name[Name.size() - 5] ==
's')
4786 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4788 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4792 Rep = Builder.CreateFMA(
A,
B,
C);
4800 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4804 if (VecWidth == 128 && EltWidth == 32)
4805 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4806 else if (VecWidth == 256 && EltWidth == 32)
4807 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4808 else if (VecWidth == 128 && EltWidth == 64)
4809 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4810 else if (VecWidth == 256 && EltWidth == 64)
4811 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4817 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4818 Rep = Builder.CreateIntrinsic(IID,
Ops);
4819 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4820 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4821 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4822 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4823 bool IsMask3 = Name[11] ==
'3';
4824 bool IsMaskZ = Name[11] ==
'z';
4826 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4827 bool IsSubAdd = Name[3] ==
's';
4831 if (Name[Name.size() - 5] ==
's')
4832 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4834 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4839 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4841 Rep = Builder.CreateIntrinsic(IID,
Ops);
4850 Value *Odd = Builder.CreateCall(FMA,
Ops);
4851 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4852 Value *Even = Builder.CreateCall(FMA,
Ops);
4858 for (
int i = 0; i != NumElts; ++i)
4859 Idxs[i] = i + (i % 2) * NumElts;
4861 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4869 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4870 Name.starts_with(
"avx512.maskz.pternlog.")) {
4871 bool ZeroMask = Name[11] ==
'z';
4875 if (VecWidth == 128 && EltWidth == 32)
4876 IID = Intrinsic::x86_avx512_pternlog_d_128;
4877 else if (VecWidth == 256 && EltWidth == 32)
4878 IID = Intrinsic::x86_avx512_pternlog_d_256;
4879 else if (VecWidth == 512 && EltWidth == 32)
4880 IID = Intrinsic::x86_avx512_pternlog_d_512;
4881 else if (VecWidth == 128 && EltWidth == 64)
4882 IID = Intrinsic::x86_avx512_pternlog_q_128;
4883 else if (VecWidth == 256 && EltWidth == 64)
4884 IID = Intrinsic::x86_avx512_pternlog_q_256;
4885 else if (VecWidth == 512 && EltWidth == 64)
4886 IID = Intrinsic::x86_avx512_pternlog_q_512;
4892 Rep = Builder.CreateIntrinsic(IID, Args);
4896 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4897 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4898 bool ZeroMask = Name[11] ==
'z';
4899 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4902 if (VecWidth == 128 && !
High)
4903 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4904 else if (VecWidth == 256 && !
High)
4905 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4906 else if (VecWidth == 512 && !
High)
4907 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4908 else if (VecWidth == 128 &&
High)
4909 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4910 else if (VecWidth == 256 &&
High)
4911 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4912 else if (VecWidth == 512 &&
High)
4913 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4919 Rep = Builder.CreateIntrinsic(IID, Args);
4923 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4924 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4925 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4926 bool ZeroMask = Name[11] ==
'z';
4927 bool IndexForm = Name[17] ==
'i';
4929 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4930 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4931 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4932 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4933 bool ZeroMask = Name[11] ==
'z';
4934 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4937 if (VecWidth == 128 && !IsSaturating)
4938 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4939 else if (VecWidth == 256 && !IsSaturating)
4940 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4941 else if (VecWidth == 512 && !IsSaturating)
4942 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4943 else if (VecWidth == 128 && IsSaturating)
4944 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4945 else if (VecWidth == 256 && IsSaturating)
4946 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4947 else if (VecWidth == 512 && IsSaturating)
4948 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4958 if (Args[1]->
getType()->isVectorTy() &&
4961 ->isIntegerTy(32) &&
4962 Args[2]->
getType()->isVectorTy() &&
4965 ->isIntegerTy(32)) {
4966 Type *NewArgType =
nullptr;
4967 if (VecWidth == 128)
4969 else if (VecWidth == 256)
4971 else if (VecWidth == 512)
4977 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4978 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4981 Rep = Builder.CreateIntrinsic(IID, Args);
4985 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4986 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4987 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4988 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4989 bool ZeroMask = Name[11] ==
'z';
4990 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4993 if (VecWidth == 128 && !IsSaturating)
4994 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4995 else if (VecWidth == 256 && !IsSaturating)
4996 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4997 else if (VecWidth == 512 && !IsSaturating)
4998 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4999 else if (VecWidth == 128 && IsSaturating)
5000 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5001 else if (VecWidth == 256 && IsSaturating)
5002 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5003 else if (VecWidth == 512 && IsSaturating)
5004 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5014 if (Args[1]->
getType()->isVectorTy() &&
5017 ->isIntegerTy(32) &&
5018 Args[2]->
getType()->isVectorTy() &&
5021 ->isIntegerTy(32)) {
5022 Type *NewArgType =
nullptr;
5023 if (VecWidth == 128)
5025 else if (VecWidth == 256)
5027 else if (VecWidth == 512)
5033 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5034 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5037 Rep = Builder.CreateIntrinsic(IID, Args);
5041 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5042 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5043 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5045 if (Name[0] ==
'a' && Name.back() ==
'2')
5046 IID = Intrinsic::x86_addcarry_32;
5047 else if (Name[0] ==
'a' && Name.back() ==
'4')
5048 IID = Intrinsic::x86_addcarry_64;
5049 else if (Name[0] ==
's' && Name.back() ==
'2')
5050 IID = Intrinsic::x86_subborrow_32;
5051 else if (Name[0] ==
's' && Name.back() ==
'4')
5052 IID = Intrinsic::x86_subborrow_64;
5059 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5062 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5065 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5069 }
else if (Name.starts_with(
"avx512.mask.") &&
5072 }
else if (Name.starts_with(
"bmi.pdep.")) {
5074 }
else if (Name.starts_with(
"bmi.pext.")) {
5084 if (Name.starts_with(
"neon.bfcvt")) {
5085 if (Name.starts_with(
"neon.bfcvtn2")) {
5087 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5089 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5090 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5093 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5094 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5096 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5100 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5101 return Builder.CreateShuffleVector(
5104 return Builder.CreateFPTrunc(CI->
getOperand(0),
5107 }
else if (Name.starts_with(
"sve.fcvt")) {
5110 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5111 .
Case(
"sve.fcvtnt.bf16f32",
5112 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5124 if (Args[1]->
getType() != BadPredTy)
5127 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5128 BadPredTy, Args[1]);
5129 Args[1] = Builder.CreateIntrinsic(
5130 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5132 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5136 if (Name ==
"neon.vcvtfp2hf")
5137 return Builder.CreateBitCast(
5138 Builder.CreateFPTrunc(
5142 if (Name ==
"neon.vcvthf2fp")
5143 return Builder.CreateFPExt(
5144 Builder.CreateBitCast(
5154 if (Name ==
"mve.vctp64.old") {
5157 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5160 Value *C1 = Builder.CreateIntrinsic(
5161 Intrinsic::arm_mve_pred_v2i,
5163 return Builder.CreateIntrinsic(
5164 Intrinsic::arm_mve_pred_i2v,
5166 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5167 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5168 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5169 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5171 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5172 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5173 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5174 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5176 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5177 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5178 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5179 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5180 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5181 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5182 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5183 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5184 std::vector<Type *> Tys;
5188 case Intrinsic::arm_mve_mull_int_predicated:
5189 case Intrinsic::arm_mve_vqdmull_predicated:
5190 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5193 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5194 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5195 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5199 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5203 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5207 case Intrinsic::arm_cde_vcx1q_predicated:
5208 case Intrinsic::arm_cde_vcx1qa_predicated:
5209 case Intrinsic::arm_cde_vcx2q_predicated:
5210 case Intrinsic::arm_cde_vcx2qa_predicated:
5211 case Intrinsic::arm_cde_vcx3q_predicated:
5212 case Intrinsic::arm_cde_vcx3qa_predicated:
5219 std::vector<Value *>
Ops;
5221 Type *Ty =
Op->getType();
5222 if (Ty->getScalarSizeInBits() == 1) {
5223 Value *C1 = Builder.CreateIntrinsic(
5224 Intrinsic::arm_mve_pred_v2i,
5226 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5231 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5246 auto UpgradeLegacyWMMAIUIntrinsicCall =
5251 Args.push_back(Builder.getFalse());
5255 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5262 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5266 NewCall->copyMetadata(*CI);
5270 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5271 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5272 "intrinsic should have 7 arguments");
5275 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5277 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5278 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5279 "intrinsic should have 8 arguments");
5284 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5287 switch (
F->getIntrinsicID()) {
5290 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5291 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5292 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5293 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5294 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5295 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5310 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5313 F->getParent(),
F->getIntrinsicID(), Overloads);
5318 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5322 NewCall->copyMetadata(*CI);
5323 NewCall->takeName(CI);
5328 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5334 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5335 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5343 if (Name.starts_with(
"addrspacecast.nonnull")) {
5346 Value *ASC = Builder.CreateAddrSpaceCast(
5369 if (NumOperands < 3)
5382 bool IsVolatile =
false;
5386 if (NumOperands > 3)
5391 if (NumOperands > 5) {
5393 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5407 if (VT->getElementType()->isIntegerTy(16)) {
5410 Val = Builder.CreateBitCast(Val, AsBF16);
5418 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5420 unsigned AddrSpace = PtrTy->getAddressSpace();
5423 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5425 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5430 MDNode *RangeNotPrivate =
5433 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5439 return Builder.CreateBitCast(RMW, RetTy);
5460 return MAV->getMetadata();
5469 if (Name ==
"label") {
5471 }
else if (Name ==
"assign") {
5478 }
else if (Name ==
"declare") {
5482 }
else if (Name ==
"addr") {
5492 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5493 }
else if (Name ==
"value") {
5496 unsigned ExprOp = 2;
5511 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5519 int64_t OffsetVal =
Offset->getSExtValue();
5520 return Builder.CreateIntrinsic(OffsetVal >= 0
5521 ? Intrinsic::vector_splice_left
5522 : Intrinsic::vector_splice_right,
5524 {CI->getArgOperand(0), CI->getArgOperand(1),
5525 Builder.getInt32(std::abs(OffsetVal))});
5530 if (Name.starts_with(
"to.fp16")) {
5532 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5533 return Builder.CreateBitCast(Cast, CI->
getType());
5536 if (Name.starts_with(
"from.fp16")) {
5538 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5539 return Builder.CreateFPExt(Cast, CI->
getType());
5598 else if (Opcode == Instruction::ICmp)
5601 else if (Opcode == Instruction::FCmp)
5604 else if (Opcode == Instruction::Select)
5609 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5621 if (Defaults.empty())
5624 unsigned OldArgCount = CI->
arg_size();
5625 unsigned NewArgCount = NewFn->
arg_size();
5627 if (OldArgCount < FirstDefault)
5631 if (OldArgCount > NewArgCount)
5636 if (OldArgCount == NewArgCount) {
5648 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5649 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5650 "missing argument outside the default range");
5651 Type *ParamTy = NewFT->getParamType(Idx);
5656 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5662 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5694 if (!Name.consume_front(
"llvm."))
5697 bool IsX86 = Name.consume_front(
"x86.");
5698 bool IsNVVM = Name.consume_front(
"nvvm.");
5699 bool IsAArch64 = Name.consume_front(
"aarch64.");
5700 bool IsARM = Name.consume_front(
"arm.");
5701 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5702 bool IsDbg = Name.consume_front(
"dbg.");
5704 (Name.consume_front(
"experimental.vector.splice") ||
5705 Name.consume_front(
"vector.splice")) &&
5706 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5707 Value *Rep =
nullptr;
5709 if (!IsX86 && Name ==
"stackprotectorcheck") {
5711 }
else if (IsNVVM) {
5715 }
else if (IsAArch64) {
5719 }
else if (IsAMDGCN) {
5723 }
else if (IsOldSplice) {
5725 }
else if (Name.consume_front(
"convert.")) {
5727 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5742 const auto &DefaultCase = [&]() ->
void {
5750 "Unknown function for CallBase upgrade and isn't just a name change");
5758 "Return type must have changed");
5759 assert(OldST->getNumElements() ==
5761 "Must have same number of elements");
5764 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5767 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5768 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5769 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5790 case Intrinsic::arm_neon_vst1:
5791 case Intrinsic::arm_neon_vst2:
5792 case Intrinsic::arm_neon_vst3:
5793 case Intrinsic::arm_neon_vst4:
5794 case Intrinsic::arm_neon_vst2lane:
5795 case Intrinsic::arm_neon_vst3lane:
5796 case Intrinsic::arm_neon_vst4lane: {
5798 NewCall = Builder.CreateCall(NewFn, Args);
5801 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5802 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5803 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5808 NewCall = Builder.CreateCall(NewFn, Args);
5811 case Intrinsic::aarch64_sve_ld3_sret:
5812 case Intrinsic::aarch64_sve_ld4_sret:
5813 case Intrinsic::aarch64_sve_ld2_sret: {
5821 Name = Name.substr(5);
5828 unsigned MinElts = RetTy->getMinNumElements() /
N;
5830 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5832 for (
unsigned I = 0;
I <
N;
I++) {
5833 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5834 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5840 case Intrinsic::coro_end_async:
5841 case Intrinsic::coro_end: {
5843 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5845 NewCall = Builder.CreateCall(NewFn, Args);
5850 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5851 Value *InRamp = Builder.CreateCall(IsInRamp);
5861 case Intrinsic::vector_extract: {
5863 Name = Name.substr(5);
5864 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5869 unsigned MinElts = RetTy->getMinNumElements();
5872 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5876 case Intrinsic::vector_insert: {
5878 Name = Name.substr(5);
5879 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5883 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5888 NewCall = Builder.CreateCall(
5892 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5898 assert(
N > 1 &&
"Create is expected to be between 2-4");
5901 unsigned MinElts = RetTy->getMinNumElements() /
N;
5902 for (
unsigned I = 0;
I <
N;
I++) {
5904 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5911 case Intrinsic::arm_neon_bfdot:
5912 case Intrinsic::arm_neon_bfmmla:
5913 case Intrinsic::arm_neon_bfmlalb:
5914 case Intrinsic::arm_neon_bfmlalt:
5915 case Intrinsic::aarch64_neon_bfdot:
5916 case Intrinsic::aarch64_neon_bfmmla:
5917 case Intrinsic::aarch64_neon_bfmlalb:
5918 case Intrinsic::aarch64_neon_bfmlalt: {
5921 "Mismatch between function args and call args");
5922 size_t OperandWidth =
5924 assert((OperandWidth == 64 || OperandWidth == 128) &&
5925 "Unexpected operand width");
5927 auto Iter = CI->
args().begin();
5928 Args.push_back(*Iter++);
5929 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5930 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5931 NewCall = Builder.CreateCall(NewFn, Args);
5935 case Intrinsic::bitreverse:
5936 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5939 case Intrinsic::ctlz:
5940 case Intrinsic::cttz: {
5947 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5951 case Intrinsic::objectsize: {
5952 Value *NullIsUnknownSize =
5956 NewCall = Builder.CreateCall(
5961 case Intrinsic::ctpop:
5962 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5964 case Intrinsic::dbg_value: {
5966 Name = Name.substr(5);
5968 if (Name.starts_with(
"dbg.addr")) {
5982 if (
Offset->isNullValue()) {
5983 NewCall = Builder.CreateCall(
5992 case Intrinsic::ptr_annotation:
6000 NewCall = Builder.CreateCall(
6009 case Intrinsic::var_annotation:
6016 NewCall = Builder.CreateCall(
6025 case Intrinsic::riscv_aes32dsi:
6026 case Intrinsic::riscv_aes32dsmi:
6027 case Intrinsic::riscv_aes32esi:
6028 case Intrinsic::riscv_aes32esmi:
6029 case Intrinsic::riscv_sm4ks:
6030 case Intrinsic::riscv_sm4ed: {
6040 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6041 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6047 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6048 Value *Res = NewCall;
6050 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6056 case Intrinsic::nvvm_mapa_shared_cluster: {
6060 Value *Res = NewCall;
6061 Res = Builder.CreateAddrSpaceCast(
6068 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
6069 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6072 Args[0] = Builder.CreateAddrSpaceCast(
6075 NewCall = Builder.CreateCall(NewFn, Args);
6082#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6083 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6085#undef G2S_CLUSTER_CASE
6090 Args[0] = Builder.CreateAddrSpaceCast(
6096 Args.push_back(Builder.getInt32(0));
6098 NewCall = Builder.CreateCall(NewFn, Args);
6105#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6106 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6114 "expected only the trailing validate_pattern to be missing");
6115 Args.push_back(Builder.getInt32(0));
6117 NewCall = Builder.CreateCall(NewFn, Args);
6123#undef NVVM_TMA_G2S_MODES
6126 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6127 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6128 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6129 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6130 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6131 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6132 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6133 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6135 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6139 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6140 NewCall = Builder.CreateCall(NewFn, Args);
6143 case Intrinsic::nvvm_tcgen05_mma_shared:
6144 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
6145 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
6146 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
6147 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
6148 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
6149 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
6150 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
6151 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
6152 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
6153 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
6154 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
6155 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
6156 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
6157 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
6158 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
6159 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
6160 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
6161 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
6162 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
6163 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
6164 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
6165 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
6166 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
6167 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
6168 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
6169 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
6170 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
6171 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
6172 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
6173 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
6174 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
6175 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
6176 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
6177 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
6178 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
6179 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
6180 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
6181 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
6183 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
6184 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
6186 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
6187 case Intrinsic::nvvm_tcgen05_mma_tensor:
6188 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
6189 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
6190 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
6191 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
6192 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
6193 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
6194 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
6195 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
6196 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
6197 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
6198 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
6199 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
6200 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
6201 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
6203 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
6204 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
6206 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
6208 Args.push_back(Builder.getInt32(0));
6209 NewCall = Builder.CreateCall(NewFn, Args);
6212 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6213 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6214 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6215 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6218 Builder.getFalse()});
6220 case Intrinsic::nvvm_mbarrier_init: {
6224 if (Args.size() == 2)
6225 Args.push_back(Builder.getInt32(0));
6226 NewCall = Builder.CreateCall(NewFn, Args);
6229 case Intrinsic::riscv_sha256sig0:
6230 case Intrinsic::riscv_sha256sig1:
6231 case Intrinsic::riscv_sha256sum0:
6232 case Intrinsic::riscv_sha256sum1:
6233 case Intrinsic::riscv_sm3p0:
6234 case Intrinsic::riscv_sm3p1: {
6241 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6243 NewCall = Builder.CreateCall(NewFn, Arg);
6245 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6252 case Intrinsic::x86_xop_vfrcz_ss:
6253 case Intrinsic::x86_xop_vfrcz_sd:
6254 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6257 case Intrinsic::x86_xop_vpermil2pd:
6258 case Intrinsic::x86_xop_vpermil2ps:
6259 case Intrinsic::x86_xop_vpermil2pd_256:
6260 case Intrinsic::x86_xop_vpermil2ps_256: {
6264 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6265 NewCall = Builder.CreateCall(NewFn, Args);
6269 case Intrinsic::x86_sse41_ptestc:
6270 case Intrinsic::x86_sse41_ptestz:
6271 case Intrinsic::x86_sse41_ptestnzc: {
6285 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6286 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6288 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6292 case Intrinsic::x86_rdtscp: {
6298 NewCall = Builder.CreateCall(NewFn);
6300 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6303 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6311 case Intrinsic::x86_sse41_insertps:
6312 case Intrinsic::x86_sse41_dppd:
6313 case Intrinsic::x86_sse41_dpps:
6314 case Intrinsic::x86_sse41_mpsadbw:
6315 case Intrinsic::x86_avx_dp_ps_256:
6316 case Intrinsic::x86_avx2_mpsadbw: {
6322 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6323 NewCall = Builder.CreateCall(NewFn, Args);
6327 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6328 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6329 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6330 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6331 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6332 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6338 NewCall = Builder.CreateCall(NewFn, Args);
6347 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6348 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6349 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6350 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6351 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6352 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6356 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6357 Args[1] = Builder.CreateBitCast(
6360 NewCall = Builder.CreateCall(NewFn, Args);
6361 Value *Res = Builder.CreateBitCast(
6369 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6370 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6371 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6375 Args[1] = Builder.CreateBitCast(
6377 Args[2] = Builder.CreateBitCast(
6380 NewCall = Builder.CreateCall(NewFn, Args);
6384 case Intrinsic::thread_pointer: {
6385 NewCall = Builder.CreateCall(NewFn, {});
6389 case Intrinsic::memcpy:
6390 case Intrinsic::memmove:
6391 case Intrinsic::memset: {
6407 NewCall = Builder.CreateCall(NewFn, Args);
6410 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6411 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6412 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6417 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6420 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6424 case Intrinsic::masked_load:
6425 case Intrinsic::masked_gather:
6426 case Intrinsic::masked_store:
6427 case Intrinsic::masked_scatter: {
6433 auto GetMaybeAlign = [](
Value *
Op) {
6435 uint64_t Val = CI->getZExtValue();
6443 auto GetAlign = [&](
Value *
Op) {
6452 case Intrinsic::masked_load:
6453 NewCall = Builder.CreateMaskedLoad(
6457 case Intrinsic::masked_gather:
6458 NewCall = Builder.CreateMaskedGather(
6464 case Intrinsic::masked_store:
6465 NewCall = Builder.CreateMaskedStore(
6469 case Intrinsic::masked_scatter:
6470 NewCall = Builder.CreateMaskedScatter(
6472 DL.getValueOrABITypeAlignment(
6486 case Intrinsic::lifetime_start:
6487 case Intrinsic::lifetime_end: {
6499 NewCall = Builder.CreateLifetimeStart(Ptr);
6501 NewCall = Builder.CreateLifetimeEnd(Ptr);
6510 case Intrinsic::x86_avx512_vpdpbusd_128:
6511 case Intrinsic::x86_avx512_vpdpbusd_256:
6512 case Intrinsic::x86_avx512_vpdpbusd_512:
6513 case Intrinsic::x86_avx512_vpdpbusds_128:
6514 case Intrinsic::x86_avx512_vpdpbusds_256:
6515 case Intrinsic::x86_avx512_vpdpbusds_512:
6516 case Intrinsic::x86_avx2_vpdpbssd_128:
6517 case Intrinsic::x86_avx2_vpdpbssd_256:
6518 case Intrinsic::x86_avx10_vpdpbssd_512:
6519 case Intrinsic::x86_avx2_vpdpbssds_128:
6520 case Intrinsic::x86_avx2_vpdpbssds_256:
6521 case Intrinsic::x86_avx10_vpdpbssds_512:
6522 case Intrinsic::x86_avx2_vpdpbsud_128:
6523 case Intrinsic::x86_avx2_vpdpbsud_256:
6524 case Intrinsic::x86_avx10_vpdpbsud_512:
6525 case Intrinsic::x86_avx2_vpdpbsuds_128:
6526 case Intrinsic::x86_avx2_vpdpbsuds_256:
6527 case Intrinsic::x86_avx10_vpdpbsuds_512:
6528 case Intrinsic::x86_avx2_vpdpbuud_128:
6529 case Intrinsic::x86_avx2_vpdpbuud_256:
6530 case Intrinsic::x86_avx10_vpdpbuud_512:
6531 case Intrinsic::x86_avx2_vpdpbuuds_128:
6532 case Intrinsic::x86_avx2_vpdpbuuds_256:
6533 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6538 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6539 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6541 NewCall = Builder.CreateCall(NewFn, Args);
6544 case Intrinsic::x86_avx512_vpdpwssd_128:
6545 case Intrinsic::x86_avx512_vpdpwssd_256:
6546 case Intrinsic::x86_avx512_vpdpwssd_512:
6547 case Intrinsic::x86_avx512_vpdpwssds_128:
6548 case Intrinsic::x86_avx512_vpdpwssds_256:
6549 case Intrinsic::x86_avx512_vpdpwssds_512:
6550 case Intrinsic::x86_avx2_vpdpwsud_128:
6551 case Intrinsic::x86_avx2_vpdpwsud_256:
6552 case Intrinsic::x86_avx10_vpdpwsud_512:
6553 case Intrinsic::x86_avx2_vpdpwsuds_128:
6554 case Intrinsic::x86_avx2_vpdpwsuds_256:
6555 case Intrinsic::x86_avx10_vpdpwsuds_512:
6556 case Intrinsic::x86_avx2_vpdpwusd_128:
6557 case Intrinsic::x86_avx2_vpdpwusd_256:
6558 case Intrinsic::x86_avx10_vpdpwusd_512:
6559 case Intrinsic::x86_avx2_vpdpwusds_128:
6560 case Intrinsic::x86_avx2_vpdpwusds_256:
6561 case Intrinsic::x86_avx10_vpdpwusds_512:
6562 case Intrinsic::x86_avx2_vpdpwuud_128:
6563 case Intrinsic::x86_avx2_vpdpwuud_256:
6564 case Intrinsic::x86_avx10_vpdpwuud_512:
6565 case Intrinsic::x86_avx2_vpdpwuuds_128:
6566 case Intrinsic::x86_avx2_vpdpwuuds_256:
6567 case Intrinsic::x86_avx10_vpdpwuuds_512:
6572 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6573 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6575 NewCall = Builder.CreateCall(NewFn, Args);
6578 assert(NewCall &&
"Should have either set this variable or returned through "
6579 "the default case");
6586 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6600 F->eraseFromParent();
6606 if (NumOperands == 0)
6614 if (NumOperands == 3) {
6618 Metadata *Elts2[] = {ScalarType, ScalarType,
6632 if (
Opc != Instruction::BitCast)
6636 Type *SrcTy = V->getType();
6653 if (
Opc != Instruction::BitCast)
6656 Type *SrcTy =
C->getType();
6673 if (Flag.getNumOperands() < 3)
6674 return std::nullopt;
6676 return Name->getString();
6677 return std::nullopt;
6691 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6692 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6693 if (auto Name = getModuleFlagNameSafely(*Flag))
6694 return *Name ==
"Debug Info Version";
6697 if (OpIt != ModFlags->op_end()) {
6698 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6705 bool BrokenDebugInfo =
false;
6708 if (!BrokenDebugInfo)
6714 M.getContext().diagnose(Diag);
6721 M.getContext().diagnose(DiagVersion);
6731 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6734 if (
F->hasFnAttribute(Attr)) {
6737 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6739 auto [Part, Rest] = S.
split(
',');
6745 const unsigned Dim = DimC -
'x';
6746 assert(Dim < 3 &&
"Unexpected dim char");
6756 F->addFnAttr(Attr, NewAttr);
6760 return S ==
"x" || S ==
"y" || S ==
"z";
6765 if (
K ==
"kernel") {
6777 const unsigned Idx = (AlignIdxValuePair >> 16);
6778 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6783 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6788 if (
K ==
"minctasm") {
6793 if (
K ==
"maxnreg") {
6798 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6802 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6806 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6810 if (
K ==
"grid_constant") {
6825 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6832 if (!SeenNodes.
insert(MD).second)
6839 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6846 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6848 const MDOperand &V = MD->getOperand(j + 1);
6854 if (NewOperands.
size() > 1)
6867 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6868 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6869 if (ModRetainReleaseMarker) {
6875 ID->getString().split(ValueComp,
"#");
6876 if (ValueComp.
size() == 2) {
6877 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6881 M.eraseNamedMetadata(ModRetainReleaseMarker);
6892 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6918 bool InvalidCast =
false;
6920 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6933 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6935 Args.push_back(Arg);
6942 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6947 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6960 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6968 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6969 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6970 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6971 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6972 {
"objc_autoreleaseReturnValue",
6973 llvm::Intrinsic::objc_autoreleaseReturnValue},
6974 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6975 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6976 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6977 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6978 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6979 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6980 {
"objc_release", llvm::Intrinsic::objc_release},
6981 {
"objc_retain", llvm::Intrinsic::objc_retain},
6982 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6983 {
"objc_retainAutoreleaseReturnValue",
6984 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6985 {
"objc_retainAutoreleasedReturnValue",
6986 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6987 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6988 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6989 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6990 {
"objc_unsafeClaimAutoreleasedReturnValue",
6991 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6992 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6993 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6994 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6995 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6996 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6997 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6998 {
"objc_arc_annotation_topdown_bbstart",
6999 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7000 {
"objc_arc_annotation_topdown_bbend",
7001 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7002 {
"objc_arc_annotation_bottomup_bbstart",
7003 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7004 {
"objc_arc_annotation_bottomup_bbend",
7005 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7007 for (
auto &
I : RuntimeFuncs)
7008 UpgradeToIntrinsic(
I.first,
I.second);
7032 std::optional<bool> UseAddressDisc;
7035 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7036 for (
const MDNode *Flag : ModFlags->operands()) {
7038 if (Name && (*Name ==
"ptrauth-init-fini" ||
7039 *Name ==
"ptrauth-init-fini-address-discrimination"))
7044 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7045 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7046 constexpr unsigned ExpectedAddressMarker = 1;
7049 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7052 bool HasAddressDisc;
7053 if (!CPA->hasAddressDiscriminator())
7054 HasAddressDisc =
false;
7055 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7056 HasAddressDisc =
true;
7060 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7063 UseAddressDisc = HasAddressDisc;
7064 return CPA->getPointer();
7068 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7071 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7073 if (!GV || !GV->hasInitializer())
7077 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7080 std::vector<Constant *> NewStructors;
7081 NewStructors.reserve(OldStructorsArray->getNumOperands());
7083 for (
Use &U : OldStructorsArray->operands()) {
7092 Func = UpgradeSinglePointer(Func);
7096 NewStructors.push_back(
7105 if (GlobalArraysToUpgrade.
empty())
7107 assert(UseAddressDisc.has_value());
7109 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7110 GV->setInitializer(NewInit);
7113 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7123 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7127 bool HasObjCFlag =
false, HasClassProperties =
false;
7128 bool HasSwiftVersionFlag =
false;
7129 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7136 if (
Op->getNumOperands() != 3)
7150 if (ID->getString() ==
"Objective-C Image Info Version")
7152 if (ID->getString() ==
"Objective-C Class Properties")
7153 HasClassProperties =
true;
7155 if (ID->getString() ==
"PIC Level") {
7156 if (
auto *Behavior =
7158 uint64_t V = Behavior->getLimitedValue();
7164 if (ID->getString() ==
"PIE Level")
7165 if (
auto *Behavior =
7172 if (ID->getString() ==
"branch-target-enforcement" ||
7173 ID->getString().starts_with(
"sign-return-address")) {
7174 if (
auto *Behavior =
7180 Op->getOperand(1),
Op->getOperand(2)};
7190 if (ID->getString() ==
"Objective-C Image Info Section") {
7193 Value->getString().split(ValueComp,
" ");
7194 if (ValueComp.
size() != 1) {
7195 std::string NewValue;
7196 for (
auto &S : ValueComp)
7197 NewValue += S.str();
7208 if (ID->getString() ==
"Objective-C Garbage Collection") {
7211 assert(Md->getValue() &&
"Expected non-empty metadata");
7212 auto Type = Md->getValue()->getType();
7215 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7216 if ((Val & 0xff) != Val) {
7217 HasSwiftVersionFlag =
true;
7218 SwiftABIVersion = (Val & 0xff00) >> 8;
7219 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7220 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7231 if (ID->getString() ==
"amdgpu_code_object_version") {
7234 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7243 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7272 if (HasObjCFlag && !HasClassProperties) {
7278 if (HasSwiftVersionFlag) {
7282 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7284 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7292 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7296 auto MatchesVersion = [](
const MDNode *
Op) {
7297 return Op->getNumOperands() >= 3 &&
7311 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7312 assert(
Op->getNumOperands() >= 2 &&
7313 "Expected at least 2 operands - name and linkage type");
7325 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7336 auto TrimSpaces = [](
StringRef Section) -> std::string {
7338 Section.split(Components,
',');
7343 for (
auto Component : Components)
7344 OS <<
',' << Component.trim();
7349 for (
auto &GV : M.globals()) {
7350 if (!GV.hasSection())
7355 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7360 GV.setSection(TrimSpaces(Section));
7376struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7377 StrictFPUpgradeVisitor() =
default;
7380 if (!
Call.isStrictFP())
7386 Call.removeFnAttr(Attribute::StrictFP);
7387 Call.addFnAttr(Attribute::NoBuiltin);
7392struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7393 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7394 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7396 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7411 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7412 StrictFPUpgradeVisitor SFPV;
7417 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7418 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7419 for (
auto &Arg :
F.args())
7421 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7423 bool AddingAttrs =
false, RemovingAttrs =
false;
7424 AttrBuilder AttrsToAdd(
F.getContext());
7429 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7430 A.isValid() &&
A.isStringAttribute()) {
7431 F.setSection(
A.getValueAsString());
7433 RemovingAttrs =
true;
7437 A.isValid() &&
A.isStringAttribute()) {
7440 AddingAttrs = RemovingAttrs =
true;
7443 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7444 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7446 RemovingAttrs =
true;
7447 if (
A.getValueAsString() ==
"true") {
7448 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7457 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7460 if (
A.getValueAsBool()) {
7461 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7467 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7468 RemovingAttrs =
true;
7475 bool HandleDenormalMode =
false;
7477 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7480 DenormalFPMath = ParsedMode;
7482 AddingAttrs = RemovingAttrs =
true;
7483 HandleDenormalMode =
true;
7487 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7491 DenormalFPMathF32 = ParsedMode;
7493 AddingAttrs = RemovingAttrs =
true;
7494 HandleDenormalMode =
true;
7498 if (HandleDenormalMode)
7499 AttrsToAdd.addDenormalFPEnvAttr(
7503 F.removeFnAttrs(AttrsToRemove);
7506 F.addFnAttrs(AttrsToAdd);
7512 if (!
F.hasFnAttribute(FnAttrName))
7513 F.addFnAttr(FnAttrName,
Value);
7520 if (!
F.hasFnAttribute(FnAttrName)) {
7522 F.addFnAttr(FnAttrName);
7524 auto A =
F.getFnAttribute(FnAttrName);
7525 if (
"false" ==
A.getValueAsString())
7526 F.removeFnAttr(FnAttrName);
7527 else if (
"true" ==
A.getValueAsString()) {
7528 F.removeFnAttr(FnAttrName);
7529 F.addFnAttr(FnAttrName);
7535 Triple T(M.getTargetTriple());
7536 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7539 uint64_t BTEValue = 0;
7540 uint64_t BPPLRValue = 0;
7541 uint64_t GCSValue = 0;
7542 uint64_t SRAValue = 0;
7543 uint64_t SRAALLValue = 0;
7544 uint64_t SRABKeyValue = 0;
7546 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7550 if (
Op->getNumOperands() != 3)
7559 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7560 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7561 : IDStr ==
"guarded-control-stack" ? &GCSValue
7562 : IDStr ==
"sign-return-address" ? &SRAValue
7563 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7564 : IDStr ==
"sign-return-address-with-bkey"
7570 *ValPtr = CI->getZExtValue();
7576 bool BTE = BTEValue == 1;
7577 bool BPPLR = BPPLRValue == 1;
7578 bool GCS = GCSValue == 1;
7579 bool SRA = SRAValue == 1;
7582 if (SRA && SRAALLValue == 1)
7583 SignTypeValue =
"all";
7586 if (SRA && SRABKeyValue == 1)
7587 SignKeyValue =
"b_key";
7589 for (
Function &
F : M.getFunctionList()) {
7590 if (
F.isDeclaration())
7597 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7598 A.isValid() &&
"none" ==
A.getValueAsString()) {
7599 F.removeFnAttr(
"sign-return-address");
7600 F.removeFnAttr(
"sign-return-address-key");
7616 if (SRAALLValue == 1)
7618 if (SRABKeyValue == 1)
7645 if (
T->getNumOperands() < 1)
7650 if (S->getString().starts_with(
"llvm.vectorizer."))
7656 StringRef OldPrefix =
"llvm.vectorizer.";
7659 if (OldTag ==
"llvm.vectorizer.unroll")
7671 if (
T->getNumOperands() < 1)
7683 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7696 Ops.reserve(
T->getNumOperands());
7697 Ops.push_back(NewTag);
7698 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7699 Ops.push_back(
T->getOperand(
I));
7716 if (
T->isDistinct()) {
7717 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7729 Ops.reserve(
T->getNumOperands());
7740 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7741 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7742 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7745 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7747 auto I =
DL.find(
"-n64-");
7749 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7754 std::string Res =
DL.str();
7757 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7758 Res.append(Res.empty() ?
"G1" :
"-G1");
7766 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7767 Res.append(
"-ni:7:8:9");
7769 if (
DL.ends_with(
"ni:7"))
7771 if (
DL.ends_with(
"ni:7:8"))
7776 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7777 Res.append(
"-p7:160:256:256:32");
7778 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7779 Res.append(
"-p8:128:128:128:48");
7780 constexpr StringRef OldP8(
"-p8:128:128-");
7781 if (
DL.contains(OldP8))
7782 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7783 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7784 Res.append(
"-p9:192:256:256:32");
7789 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7790 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7791 Res.append((
"-" + AS +
":32:32").str());
7796 if (!
DL.contains(
"m:e"))
7797 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7802 if (
T.isSystemZ() && !
DL.empty()) {
7804 if (!
DL.contains(
"-S64"))
7805 return "E-S64" +
DL.drop_front(1).str();
7809 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7812 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7813 if (!
DL.contains(AddrSpaces)) {
7815 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7816 if (R.match(Res, &
Groups))
7822 if (
T.isAArch64()) {
7824 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7825 Res.append(
"-Fn32");
7826 AddPtr32Ptr64AddrSpaces();
7830 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7834 std::string I64 =
"-i64:64";
7835 std::string I128 =
"-i128:128";
7837 size_t Pos = Res.find(I64);
7838 if (Pos !=
size_t(-1))
7839 Res.insert(Pos + I64.size(), I128);
7843 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7844 size_t Pos = Res.find(
"-S128");
7847 Res.insert(Pos,
"-f64:32:64");
7853 AddPtr32Ptr64AddrSpaces();
7861 if (!
T.isOSIAMCU()) {
7862 std::string I128 =
"-i128:128";
7865 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7866 if (R.match(Res, &
Groups))
7874 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7876 auto I =
Ref.find(
"-f80:32-");
7878 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7886 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7889 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7890 B.removeAttribute(
"no-frame-pointer-elim");
7892 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7894 if (FramePointer !=
"all")
7895 FramePointer =
"non-leaf";
7896 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7898 if (!FramePointer.
empty())
7899 B.addAttribute(
"frame-pointer", FramePointer);
7901 A =
B.getAttribute(
"null-pointer-is-valid");
7904 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7905 B.removeAttribute(
"null-pointer-is-valid");
7906 if (NullPointerIsValid)
7907 B.addAttribute(Attribute::NullPointerIsValid);
7910 A =
B.getAttribute(
"uniform-work-group-size");
7914 bool IsTrue = Val ==
"true";
7915 B.removeAttribute(
"uniform-work-group-size");
7917 B.addAttribute(
"uniform-work-group-size");
7928 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFAddUpgrade(StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static 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 unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID)
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 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.