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());
993 .
Case(
"smax", Intrinsic::smax)
994 .
Case(
"smin", Intrinsic::smin)
995 .
Case(
"umax", Intrinsic::umax)
996 .
Case(
"umin", Intrinsic::umin)
999 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1002 F->getReturnType());
1006 if (Name.starts_with(
"addp")) {
1008 if (
F->arg_size() != 2)
1011 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1013 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1019 if (Name.starts_with(
"bfcvt")) {
1025 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1032 if (Name.consume_front(
"sve.")) {
1034 if (Name.consume_front(
"bf")) {
1035 if (Name ==
"mmla") {
1036 Type *Tys[] = {
F->getReturnType(),
1037 std::next(
F->arg_begin())->getType()};
1039 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1042 if (Name.consume_back(
".lane")) {
1046 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1047 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1048 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1060 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1065 if (Name.consume_front(
"convert.from.svbool")) {
1068 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1071 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1076 if (Name.consume_front(
"convert.to.svbool")) {
1079 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1082 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1087 if (Name.consume_front(
"addqv")) {
1089 if (!
F->getReturnType()->isFPOrFPVectorTy())
1092 auto Args =
F->getFunctionType()->params();
1093 Type *Tys[] = {
F->getReturnType(), Args[1]};
1095 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1099 if (Name.consume_front(
"ld")) {
1101 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1102 if (LdRegex.
match(Name)) {
1108 "Expected 2 arguments for ld* intrinsic.");
1109 Type *PtrTy =
F->getArg(1)->getType();
1112 Intrinsic::aarch64_sve_ld2_sret,
1113 Intrinsic::aarch64_sve_ld3_sret,
1114 Intrinsic::aarch64_sve_ld4_sret,
1117 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1123 if (Name.consume_front(
"tuple.")) {
1125 if (Name.starts_with(
"get")) {
1127 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1129 F->getParent(), Intrinsic::vector_extract, Tys);
1133 if (Name.starts_with(
"set")) {
1135 auto Args =
F->getFunctionType()->params();
1136 Type *Tys[] = {Args[0], Args[2], Args[1]};
1138 F->getParent(), Intrinsic::vector_insert, Tys);
1142 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1143 if (CreateTupleRegex.
match(Name)) {
1145 auto Args =
F->getFunctionType()->params();
1146 Type *Tys[] = {
F->getReturnType(), Args[1]};
1148 F->getParent(), Intrinsic::vector_insert, Tys);
1154 if (Name.starts_with(
"rev.nxv")) {
1157 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1163 if (Name.consume_front(
"sme.")) {
1165 if (Name.consume_front(
"ftmopa.")) {
1170 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1171 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1189#define NVVM_TMA_G2S_MODES(M) \
1190 M(tile_1d, "tile.1d") \
1191 M(tile_2d, "tile.2d") \
1192 M(tile_3d, "tile.3d") \
1193 M(tile_4d, "tile.4d") \
1194 M(tile_5d, "tile.5d") \
1195 M(tile_gather4_2d, "tile.gather4.2d") \
1196 M(im2col_3d, "im2col.3d") \
1197 M(im2col_4d, "im2col.4d") \
1198 M(im2col_5d, "im2col.5d") \
1199 M(im2col_w_3d, "im2col.w.3d") \
1200 M(im2col_w_4d, "im2col.w.4d") \
1201 M(im2col_w_5d, "im2col.w.5d") \
1202 M(im2col_w_128_3d, "im2col.w.128.3d") \
1203 M(im2col_w_128_4d, "im2col.w.128.4d") \
1204 M(im2col_w_128_5d, "im2col.w.128.5d")
1216 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1219#define G2S_ID(ID_SUFFIX, NAME) \
1220 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1230 size_t NumParams =
F->getFunctionType()->getNumParams();
1234 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1241 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1242 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1243 "expected the i64 cache-hint after the multicast mask");
1244 Type *MaskTy = Params[MaskIdx];
1259 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1262#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1263 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1275 if (!
F->getFunctionType()
1276 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1292 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1297 size_t NumParams =
F->getFunctionType()->getNumParams();
1298 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1302 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1307 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1320 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1325 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1328 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1348 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1351 auto [RedOpName, ShapeName] = Name.split(
'.');
1356 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1357 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1358 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1359 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1360 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1361 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1362 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1363 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1369 if (Name.consume_front(
"mapa.shared.cluster"))
1370 if (
F->getReturnType()->getPointerAddressSpace() ==
1372 return Intrinsic::nvvm_mapa_shared_cluster;
1374 if (Name.consume_front(
"cp.async.bulk.")) {
1377 .
Case(
"shared.cta.to.cluster",
1378 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1382 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1392 if (!Name.consume_front(
"tcgen05.commit."))
1395 if (Name.consume_front(
"shared."))
1397 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1398 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1401 if (Name.consume_front(
"mc.shared.")) {
1403 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1407 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1408 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1417 if (
F->arg_size() != 2)
1420 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1421 Name.consume_front(
"tcgen05.alloc."))
1423 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1424 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1427 if (Name.consume_front(
"tcgen05.dealloc."))
1429 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1430 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1437 if (Name.consume_front(
"fma.rn."))
1439 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1440 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1441 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1442 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1445 if (Name.consume_front(
"fmax."))
1447 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1448 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1449 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1450 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1451 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1452 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1453 .
Case(
"ftz.nan.xorsign.abs.bf16",
1454 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1455 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1456 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1457 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1458 .
Case(
"ftz.xorsign.abs.bf16x2",
1459 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1460 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1461 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1462 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1463 .
Case(
"nan.xorsign.abs.bf16x2",
1464 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1465 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1466 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1469 if (Name.consume_front(
"fmin."))
1471 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1472 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1473 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1474 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1475 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1476 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1477 .
Case(
"ftz.nan.xorsign.abs.bf16",
1478 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1479 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1480 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1481 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1482 .
Case(
"ftz.xorsign.abs.bf16x2",
1483 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1484 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1485 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1486 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1487 .
Case(
"nan.xorsign.abs.bf16x2",
1488 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1489 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1490 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1493 if (Name.consume_front(
"neg."))
1495 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1496 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1505 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1510 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1513 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1516 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1517 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1523static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1525 auto [Modifiers,
Type] = Name.rsplit(
'.');
1527 return std::nullopt;
1537 return std::nullopt;
1540 .
Case(
"", Intrinsic::nvvm_fadd)
1541 .
Case(
".ftz", Intrinsic::nvvm_fadd_ftz)
1542 .
Case(
".sat", Intrinsic::nvvm_fadd_sat)
1543 .
Case(
".ftz.sat", Intrinsic::nvvm_fadd_ftz_sat)
1546 return std::nullopt;
1552 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1555 return Intrinsic::nvvm_mbarrier_init;
1559 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1560 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1561 Name.consume_front(
"param");
1565 if (!Name.consume_front(
"vp."))
1594 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1595 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1602 if (!Name.consume_front(
"vp."))
1622 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1623 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1628 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1640 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1651 if (Name.starts_with(
"to.fp16")) {
1655 FuncTy->getReturnType());
1658 if (Name.starts_with(
"from.fp16")) {
1662 FuncTy->getReturnType());
1672 if (Defaults.empty())
1675 unsigned FullArgCount = FirstDefault + Defaults.size();
1678 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1681 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1683 NumMissingTrailingParams))
1686 return FullArgCount;
1693 unsigned FullArgCount =
1695 if (FullArgCount == 0)
1701 "total number of default args does not match intrinsic signature");
1706 bool CanUpgradeDebugIntrinsicsToRecords) {
1707 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1712 if (!Name.consume_front(
"llvm.") || Name.empty())
1718 bool IsArm = Name.consume_front(
"arm.");
1719 if (IsArm || Name.consume_front(
"aarch64.")) {
1725 if (Name.consume_front(
"amdgcn.")) {
1726 if (Name ==
"alignbit") {
1729 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1733 if (Name.consume_front(
"atomic.")) {
1734 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1735 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1744 if (Name.starts_with(
"addrspacecast.nonnull")) {
1751 switch (
F->getIntrinsicID()) {
1755 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1756 if (
F->arg_size() == 7) {
1761 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1762 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1763 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1764 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1765 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1766 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1767 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1768 if (
F->arg_size() == 8) {
1775 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1776 Name.consume_front(
"flat.atomic.")) {
1777 if (Name.starts_with(
"fadd") ||
1779 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1780 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1788 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1793 if (Name.starts_with(
"ldexp.")) {
1796 F->getParent(), Intrinsic::ldexp,
1797 {F->getReturnType(), F->getArg(1)->getType()});
1806 if (
F->arg_size() == 1) {
1807 if (Name.consume_front(
"convert.")) {
1821 F->arg_begin()->getType());
1827 if (Name ==
"coro.end" &&
1828 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1829 CoroEndID = Intrinsic::coro_end;
1830 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1831 CoroEndID = Intrinsic::coro_end_async;
1842 if (Name.consume_front(
"dbg.")) {
1844 if (CanUpgradeDebugIntrinsicsToRecords) {
1845 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1846 Name ==
"declare" || Name ==
"label") {
1855 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1858 Intrinsic::dbg_value);
1865 if (Name.consume_front(
"experimental.vector.")) {
1871 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1872 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1873 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1874 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1875 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1877 Intrinsic::vector_partial_reduce_add)
1880 const auto *FT =
F->getFunctionType();
1882 if (ID == Intrinsic::vector_extract ||
1883 ID == Intrinsic::vector_interleave2)
1886 if (ID != Intrinsic::vector_interleave2)
1888 if (ID == Intrinsic::vector_insert ||
1889 ID == Intrinsic::vector_partial_reduce_add)
1897 if (Name.consume_front(
"reduce.")) {
1899 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1900 if (R.match(Name, &
Groups))
1902 .
Case(
"add", Intrinsic::vector_reduce_add)
1903 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1904 .
Case(
"and", Intrinsic::vector_reduce_and)
1905 .
Case(
"or", Intrinsic::vector_reduce_or)
1906 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1907 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1908 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1909 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1910 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1911 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1912 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1917 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1922 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1923 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1928 auto Args =
F->getFunctionType()->params();
1930 {Args[V2 ? 1 : 0]});
1936 if (Name.consume_front(
"splice"))
1940 if (Name.consume_front(
"experimental.stepvector.")) {
1944 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1949 if (Name.starts_with(
"flt.rounds")) {
1952 Intrinsic::get_rounding);
1957 if (Name.starts_with(
"invariant.group.barrier")) {
1959 auto Args =
F->getFunctionType()->params();
1960 Type* ObjectPtr[1] = {Args[0]};
1963 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1968 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1969 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1970 if (IsLifetimeStart || IsLifetimeEnd) {
1971 if (
F->arg_size() == 2) {
1972 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1973 : Intrinsic::lifetime_end;
1978 F->getArg(1)->getType());
1980 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2000 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2001 .StartsWith(
"memmove.", Intrinsic::memmove)
2003 if (
F->arg_size() == 5) {
2007 F->getFunctionType()->params().slice(0, 3);
2013 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2016 const auto *FT =
F->getFunctionType();
2017 Type *ParamTypes[2] = {
2018 FT->getParamType(0),
2022 Intrinsic::memset, ParamTypes);
2028 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2029 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2030 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2031 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2033 if (MaskedID &&
F->arg_size() == 4) {
2035 if (MaskedID == Intrinsic::masked_load ||
2036 MaskedID == Intrinsic::masked_gather) {
2038 F->getParent(), MaskedID,
2039 {F->getReturnType(), F->getArg(0)->getType()});
2043 F->getParent(), MaskedID,
2044 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2050 if (Name.consume_front(
"nvvm.")) {
2052 if (
F->arg_size() == 1) {
2055 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2056 .Case(
"clz.i", Intrinsic::ctlz)
2057 .
Case(
"popc.i", Intrinsic::ctpop)
2061 {F->getReturnType()});
2064 }
else if (
F->arg_size() == 2) {
2067 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2068 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2069 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2070 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2074 {F->getReturnType()});
2111 F->getParent(), IID,
F->getReturnType(),
2112 F->getFunctionType()->params());
2123 {F->getArg(0)->getType()});
2171 F->getArg(0)->getType());
2179 bool Expand =
false;
2180 if (Name.consume_front(
"abs."))
2183 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2184 else if (Name.consume_front(
"fabs."))
2186 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2187 else if (Name.consume_front(
"add."))
2190 else if (Name.consume_front(
"ex2.approx."))
2193 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2194 else if (Name.consume_front(
"atomic.load."))
2203 else if (Name.consume_front(
"atomic."))
2218 else if (Name.consume_front(
"bitcast."))
2221 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2222 else if (Name.consume_front(
"rotate."))
2224 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2225 else if (Name.consume_front(
"ptr.gen.to."))
2228 else if (Name.consume_front(
"ptr."))
2231 else if (Name.consume_front(
"ldg.global."))
2233 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2234 Name.starts_with(
"p."));
2237 .
Case(
"barrier0",
true)
2238 .
Case(
"barrier.n",
true)
2239 .
Case(
"barrier.sync.cnt",
true)
2240 .
Case(
"barrier.sync",
true)
2241 .
Case(
"barrier",
true)
2242 .
Case(
"bar.sync",
true)
2243 .
Case(
"barrier0.popc",
true)
2244 .
Case(
"barrier0.and",
true)
2245 .
Case(
"barrier0.or",
true)
2246 .
Case(
"clz.ll",
true)
2247 .
Case(
"popc.ll",
true)
2249 .
Case(
"swap.lo.hi.b64",
true)
2250 .
Case(
"tanh.approx.f32",
true)
2262 if (Name.starts_with(
"objectsize.")) {
2263 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2264 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2267 Intrinsic::objectsize, Tys);
2274 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2277 F->getParent(), Intrinsic::ptr_annotation,
2278 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2284 if (Name.consume_front(
"riscv.")) {
2287 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2288 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2289 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2290 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2293 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2306 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2307 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2316 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2317 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2318 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2319 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2324 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2333 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2335 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2344 if (Name ==
"stackprotectorcheck") {
2348 if (Name.starts_with(
"strip.invariant.group")) {
2353 F->getParent(), Intrinsic::launder_invariant_group,
2354 F->getReturnType());
2360 if (Name ==
"thread.pointer") {
2362 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2368 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2371 F->getParent(), Intrinsic::var_annotation,
2372 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2375 if (Name.consume_front(
"vector.splice")) {
2376 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2386 if (Name.consume_front(
"wasm.")) {
2389 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2390 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2391 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2396 F->getReturnType());
2400 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2402 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2404 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2423 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2433 std::string
Name =
F->getName().str();
2436 Name,
F->getParent());
2447 if (Result != std::nullopt) {
2464 bool CanUpgradeDebugIntrinsicsToRecords) {
2484 GV->
getName() ==
"llvm.global_dtors")) ||
2499 unsigned N =
Init->getNumOperands();
2500 std::vector<Constant *> NewCtors(
N);
2501 for (
unsigned i = 0; i !=
N; ++i) {
2504 Ctor->getAggregateElement(1),
2518 unsigned NumElts = ResultTy->getNumElements() * 8;
2522 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2532 for (
unsigned l = 0; l != NumElts; l += 16)
2533 for (
unsigned i = 0; i != 16; ++i) {
2534 unsigned Idx = NumElts + i - Shift;
2536 Idx -= NumElts - 16;
2537 Idxs[l + i] = Idx + l;
2540 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2544 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2552 unsigned NumElts = ResultTy->getNumElements() * 8;
2556 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2566 for (
unsigned l = 0; l != NumElts; l += 16)
2567 for (
unsigned i = 0; i != 16; ++i) {
2568 unsigned Idx = i + Shift;
2570 Idx += NumElts - 16;
2571 Idxs[l + i] = Idx + l;
2574 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2578 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2586 Mask = Builder.CreateBitCast(Mask, MaskTy);
2592 for (
unsigned i = 0; i != NumElts; ++i)
2594 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2605 if (
C->isAllOnesValue())
2610 return Builder.CreateSelect(Mask, Op0, Op1);
2617 if (
C->isAllOnesValue())
2621 Mask->getType()->getIntegerBitWidth());
2622 Mask = Builder.CreateBitCast(Mask, MaskTy);
2623 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2624 return Builder.CreateSelect(Mask, Op0, Op1);
2637 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2638 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2643 ShiftVal &= (NumElts - 1);
2652 if (ShiftVal > 16) {
2660 for (
unsigned l = 0; l < NumElts; l += 16) {
2661 for (
unsigned i = 0; i != 16; ++i) {
2662 unsigned Idx = ShiftVal + i;
2663 if (!IsVALIGN && Idx >= 16)
2664 Idx += NumElts - 16;
2665 Indices[l + i] = Idx + l;
2670 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2676 bool ZeroMask,
bool IndexForm) {
2679 unsigned EltWidth = Ty->getScalarSizeInBits();
2680 bool IsFloat = Ty->isFPOrFPVectorTy();
2682 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2683 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2684 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2685 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2686 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2687 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2688 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2689 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2690 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2691 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2692 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2693 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2694 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2695 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2696 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2697 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2698 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2699 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2700 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2701 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2702 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2703 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2704 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2705 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2706 else if (VecWidth == 128 && EltWidth == 16)
2707 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2708 else if (VecWidth == 256 && EltWidth == 16)
2709 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2710 else if (VecWidth == 512 && EltWidth == 16)
2711 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2712 else if (VecWidth == 128 && EltWidth == 8)
2713 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2714 else if (VecWidth == 256 && EltWidth == 8)
2715 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2716 else if (VecWidth == 512 && EltWidth == 8)
2717 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2728 Value *V = Builder.CreateIntrinsic(IID, Args);
2740 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2751 bool IsRotateRight) {
2761 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2762 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2765 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2766 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2811 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2816 bool IsShiftRight,
bool ZeroMask) {
2830 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2831 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2834 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2835 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2850 const Align Alignment =
2852 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2857 if (
C->isAllOnesValue())
2858 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2863 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2869 const Align Alignment =
2878 if (
C->isAllOnesValue())
2879 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2884 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2890 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2891 {Op0, Builder.getInt1(
false)});
2906 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2907 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2908 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2909 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2910 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2913 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2914 LHS = Builder.CreateAnd(
LHS, Mask);
2915 RHS = Builder.CreateAnd(
RHS, Mask);
2932 if (!
C || !
C->isAllOnesValue())
2933 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2938 for (
unsigned i = 0; i != NumElts; ++i)
2940 for (
unsigned i = NumElts; i != 8; ++i)
2941 Indices[i] = NumElts + i % NumElts;
2942 Vec = Builder.CreateShuffleVector(Vec,
2946 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2950 unsigned CC,
bool Signed) {
2958 }
else if (CC == 7) {
2994 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2995 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2997 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2998 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3007 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3013 Name = Name.substr(12);
3018 if (Name.starts_with(
"max.p")) {
3019 if (VecWidth == 128 && EltWidth == 32)
3020 IID = Intrinsic::x86_sse_max_ps;
3021 else if (VecWidth == 128 && EltWidth == 64)
3022 IID = Intrinsic::x86_sse2_max_pd;
3023 else if (VecWidth == 256 && EltWidth == 32)
3024 IID = Intrinsic::x86_avx_max_ps_256;
3025 else if (VecWidth == 256 && EltWidth == 64)
3026 IID = Intrinsic::x86_avx_max_pd_256;
3029 }
else if (Name.starts_with(
"min.p")) {
3030 if (VecWidth == 128 && EltWidth == 32)
3031 IID = Intrinsic::x86_sse_min_ps;
3032 else if (VecWidth == 128 && EltWidth == 64)
3033 IID = Intrinsic::x86_sse2_min_pd;
3034 else if (VecWidth == 256 && EltWidth == 32)
3035 IID = Intrinsic::x86_avx_min_ps_256;
3036 else if (VecWidth == 256 && EltWidth == 64)
3037 IID = Intrinsic::x86_avx_min_pd_256;
3040 }
else if (Name.starts_with(
"pshuf.b.")) {
3041 if (VecWidth == 128)
3042 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3043 else if (VecWidth == 256)
3044 IID = Intrinsic::x86_avx2_pshuf_b;
3045 else if (VecWidth == 512)
3046 IID = Intrinsic::x86_avx512_pshuf_b_512;
3049 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3050 if (VecWidth == 128)
3051 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3052 else if (VecWidth == 256)
3053 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3054 else if (VecWidth == 512)
3055 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3058 }
else if (Name.starts_with(
"pmulh.w.")) {
3059 if (VecWidth == 128)
3060 IID = Intrinsic::x86_sse2_pmulh_w;
3061 else if (VecWidth == 256)
3062 IID = Intrinsic::x86_avx2_pmulh_w;
3063 else if (VecWidth == 512)
3064 IID = Intrinsic::x86_avx512_pmulh_w_512;
3067 }
else if (Name.starts_with(
"pmulhu.w.")) {
3068 if (VecWidth == 128)
3069 IID = Intrinsic::x86_sse2_pmulhu_w;
3070 else if (VecWidth == 256)
3071 IID = Intrinsic::x86_avx2_pmulhu_w;
3072 else if (VecWidth == 512)
3073 IID = Intrinsic::x86_avx512_pmulhu_w_512;
3076 }
else if (Name.starts_with(
"pmaddw.d.")) {
3077 if (VecWidth == 128)
3078 IID = Intrinsic::x86_sse2_pmadd_wd;
3079 else if (VecWidth == 256)
3080 IID = Intrinsic::x86_avx2_pmadd_wd;
3081 else if (VecWidth == 512)
3082 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3085 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3086 if (VecWidth == 128)
3087 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3088 else if (VecWidth == 256)
3089 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3090 else if (VecWidth == 512)
3091 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3094 }
else if (Name.starts_with(
"packsswb.")) {
3095 if (VecWidth == 128)
3096 IID = Intrinsic::x86_sse2_packsswb_128;
3097 else if (VecWidth == 256)
3098 IID = Intrinsic::x86_avx2_packsswb;
3099 else if (VecWidth == 512)
3100 IID = Intrinsic::x86_avx512_packsswb_512;
3103 }
else if (Name.starts_with(
"packssdw.")) {
3104 if (VecWidth == 128)
3105 IID = Intrinsic::x86_sse2_packssdw_128;
3106 else if (VecWidth == 256)
3107 IID = Intrinsic::x86_avx2_packssdw;
3108 else if (VecWidth == 512)
3109 IID = Intrinsic::x86_avx512_packssdw_512;
3112 }
else if (Name.starts_with(
"packuswb.")) {
3113 if (VecWidth == 128)
3114 IID = Intrinsic::x86_sse2_packuswb_128;
3115 else if (VecWidth == 256)
3116 IID = Intrinsic::x86_avx2_packuswb;
3117 else if (VecWidth == 512)
3118 IID = Intrinsic::x86_avx512_packuswb_512;
3121 }
else if (Name.starts_with(
"packusdw.")) {
3122 if (VecWidth == 128)
3123 IID = Intrinsic::x86_sse41_packusdw;
3124 else if (VecWidth == 256)
3125 IID = Intrinsic::x86_avx2_packusdw;
3126 else if (VecWidth == 512)
3127 IID = Intrinsic::x86_avx512_packusdw_512;
3130 }
else if (Name.starts_with(
"vpermilvar.")) {
3131 if (VecWidth == 128 && EltWidth == 32)
3132 IID = Intrinsic::x86_avx_vpermilvar_ps;
3133 else if (VecWidth == 128 && EltWidth == 64)
3134 IID = Intrinsic::x86_avx_vpermilvar_pd;
3135 else if (VecWidth == 256 && EltWidth == 32)
3136 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3137 else if (VecWidth == 256 && EltWidth == 64)
3138 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3139 else if (VecWidth == 512 && EltWidth == 32)
3140 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3141 else if (VecWidth == 512 && EltWidth == 64)
3142 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3145 }
else if (Name ==
"cvtpd2dq.256") {
3146 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3147 }
else if (Name ==
"cvtpd2ps.256") {
3148 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3149 }
else if (Name ==
"cvttpd2dq.256") {
3150 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3151 }
else if (Name ==
"cvttps2dq.128") {
3152 IID = Intrinsic::x86_sse2_cvttps2dq;
3153 }
else if (Name ==
"cvttps2dq.256") {
3154 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3155 }
else if (Name.starts_with(
"permvar.")) {
3157 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3158 IID = Intrinsic::x86_avx2_permps;
3159 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3160 IID = Intrinsic::x86_avx2_permd;
3161 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3162 IID = Intrinsic::x86_avx512_permvar_df_256;
3163 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3164 IID = Intrinsic::x86_avx512_permvar_di_256;
3165 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3166 IID = Intrinsic::x86_avx512_permvar_sf_512;
3167 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3168 IID = Intrinsic::x86_avx512_permvar_si_512;
3169 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3170 IID = Intrinsic::x86_avx512_permvar_df_512;
3171 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3172 IID = Intrinsic::x86_avx512_permvar_di_512;
3173 else if (VecWidth == 128 && EltWidth == 16)
3174 IID = Intrinsic::x86_avx512_permvar_hi_128;
3175 else if (VecWidth == 256 && EltWidth == 16)
3176 IID = Intrinsic::x86_avx512_permvar_hi_256;
3177 else if (VecWidth == 512 && EltWidth == 16)
3178 IID = Intrinsic::x86_avx512_permvar_hi_512;
3179 else if (VecWidth == 128 && EltWidth == 8)
3180 IID = Intrinsic::x86_avx512_permvar_qi_128;
3181 else if (VecWidth == 256 && EltWidth == 8)
3182 IID = Intrinsic::x86_avx512_permvar_qi_256;
3183 else if (VecWidth == 512 && EltWidth == 8)
3184 IID = Intrinsic::x86_avx512_permvar_qi_512;
3187 }
else if (Name.starts_with(
"dbpsadbw.")) {
3188 if (VecWidth == 128)
3189 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3190 else if (VecWidth == 256)
3191 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3192 else if (VecWidth == 512)
3193 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3196 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3197 if (VecWidth == 128)
3198 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3199 else if (VecWidth == 256)
3200 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3201 else if (VecWidth == 512)
3202 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3205 }
else if (Name.starts_with(
"conflict.")) {
3206 if (Name[9] ==
'd' && VecWidth == 128)
3207 IID = Intrinsic::x86_avx512_conflict_d_128;
3208 else if (Name[9] ==
'd' && VecWidth == 256)
3209 IID = Intrinsic::x86_avx512_conflict_d_256;
3210 else if (Name[9] ==
'd' && VecWidth == 512)
3211 IID = Intrinsic::x86_avx512_conflict_d_512;
3212 else if (Name[9] ==
'q' && VecWidth == 128)
3213 IID = Intrinsic::x86_avx512_conflict_q_128;
3214 else if (Name[9] ==
'q' && VecWidth == 256)
3215 IID = Intrinsic::x86_avx512_conflict_q_256;
3216 else if (Name[9] ==
'q' && VecWidth == 512)
3217 IID = Intrinsic::x86_avx512_conflict_q_512;
3220 }
else if (Name.starts_with(
"pavg.")) {
3221 if (Name[5] ==
'b' && VecWidth == 128)
3222 IID = Intrinsic::x86_sse2_pavg_b;
3223 else if (Name[5] ==
'b' && VecWidth == 256)
3224 IID = Intrinsic::x86_avx2_pavg_b;
3225 else if (Name[5] ==
'b' && VecWidth == 512)
3226 IID = Intrinsic::x86_avx512_pavg_b_512;
3227 else if (Name[5] ==
'w' && VecWidth == 128)
3228 IID = Intrinsic::x86_sse2_pavg_w;
3229 else if (Name[5] ==
'w' && VecWidth == 256)
3230 IID = Intrinsic::x86_avx2_pavg_w;
3231 else if (Name[5] ==
'w' && VecWidth == 512)
3232 IID = Intrinsic::x86_avx512_pavg_w_512;
3241 Rep = Builder.CreateIntrinsic(IID, Args);
3252 if (AsmStr->find(
"mov\tfp") == 0 &&
3253 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3254 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3255 AsmStr->replace(Pos, 1,
";");
3261 Value *Rep =
nullptr;
3263 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3265 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3266 {Arg, Builder.getTrue()},
3268 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3269 Type *Ty = (Name ==
"abs.bf16")
3273 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3274 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3275 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3276 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3277 : Intrinsic::nvvm_fabs;
3278 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3279 }
else if (Name.consume_front(
"add.")) {
3282 assert(
FAdd &&
"unsupported nvvm.add.* intrinsic");
3285 Rep = Builder.CreateIntrinsic(
3287 {A, CI->getArgOperand(1),
3288 Builder.getInt32(static_cast<int>(RoundingMode))});
3289 }
else if (Name.consume_front(
"ex2.approx.")) {
3291 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3292 : Intrinsic::nvvm_ex2_approx;
3293 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3294 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3295 Name.starts_with(
"atomic.load.add.f64.p")) {
3298 Rep = Builder.CreateAtomicRMW(
3304 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3305 Name.starts_with(
"atomic.load.dec.32.p")) {
3310 Rep = Builder.CreateAtomicRMW(
3314 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3320 Op.contains(
".cta.") ?
"block" :
"");
3321 if (
Op.starts_with(
"cas.")) {
3323 Value *Pair = Builder.CreateAtomicCmpXchg(
3326 Rep = Builder.CreateExtractValue(Pair, 0);
3344 "unexpected nvvm scoped atomic intrinsic");
3345 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3348 }
else if (Name ==
"clz.ll") {
3351 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3352 {Arg, Builder.getFalse()},
3354 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3355 }
else if (Name ==
"popc.ll") {
3359 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3360 Arg,
nullptr,
"ctpop");
3361 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3362 }
else if (Name ==
"h2f") {
3364 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3365 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3366 }
else if (Name.consume_front(
"bitcast.") &&
3367 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3370 }
else if (Name ==
"rotate.b32") {
3373 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3374 {Arg, Arg, ShiftAmt});
3375 }
else if (Name ==
"rotate.b64") {
3379 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3380 {Arg, Arg, ZExtShiftAmt});
3381 }
else if (Name ==
"rotate.right.b64") {
3385 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3386 {Arg, Arg, ZExtShiftAmt});
3387 }
else if (Name ==
"swap.lo.hi.b64") {
3390 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3391 {Arg, Arg, Builder.getInt64(32)});
3392 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3395 Name.starts_with(
".to.gen"))) {
3397 }
else if (Name.consume_front(
"ldg.global")) {
3401 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3404 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3406 }
else if (Name ==
"tanh.approx.f32") {
3410 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3412 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3414 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3415 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3417 }
else if (Name ==
"barrier") {
3418 Rep = Builder.CreateIntrinsic(
3419 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3421 }
else if (Name ==
"barrier.sync") {
3422 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3424 }
else if (Name ==
"barrier.sync.cnt") {
3425 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3427 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3428 Name ==
"barrier0.or") {
3430 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3434 .
Case(
"barrier0.popc",
3435 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3436 .
Case(
"barrier0.and",
3437 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3438 .
Case(
"barrier0.or",
3439 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3440 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3441 Rep = Builder.CreateZExt(Bar, CI->
getType());
3455 ? Builder.CreateBitCast(Arg, NewType)
3458 Rep = Builder.CreateCall(NewFn, Args);
3459 if (
F->getReturnType()->isIntegerTy())
3460 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3470 Value *Rep =
nullptr;
3472 if (Name.starts_with(
"sse4a.movnt.")) {
3484 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3487 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3488 }
else if (Name.starts_with(
"avx.movnt.") ||
3489 Name.starts_with(
"avx512.storent.")) {
3501 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3502 }
else if (Name ==
"sse2.storel.dq") {
3507 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3508 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3509 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3510 }
else if (Name.starts_with(
"sse.storeu.") ||
3511 Name.starts_with(
"sse2.storeu.") ||
3512 Name.starts_with(
"avx.storeu.")) {
3515 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3516 }
else if (Name ==
"avx512.mask.store.ss") {
3520 }
else if (Name.starts_with(
"avx512.mask.store")) {
3522 bool Aligned = Name[17] !=
'u';
3525 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3528 bool CmpEq = Name[9] ==
'e';
3531 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3532 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3539 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3540 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3542 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3543 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3544 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3545 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3546 Name.starts_with(
"sse2.sqrt.p") ||
3547 Name.starts_with(
"sse.sqrt.p")) {
3548 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3549 {CI->getArgOperand(0)});
3550 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3554 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3555 : Intrinsic::x86_avx512_sqrt_pd_512;
3558 Rep = Builder.CreateIntrinsic(IID, Args);
3560 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3561 {CI->getArgOperand(0)});
3565 }
else if (Name.starts_with(
"avx512.ptestm") ||
3566 Name.starts_with(
"avx512.ptestnm")) {
3570 Rep = Builder.CreateAnd(Op0, Op1);
3576 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3578 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3581 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3584 }
else if (Name.starts_with(
"avx512.kunpck")) {
3589 for (
unsigned i = 0; i != NumElts; ++i)
3598 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3599 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3600 }
else if (Name ==
"avx512.kand.w") {
3603 Rep = Builder.CreateAnd(
LHS,
RHS);
3604 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3605 }
else if (Name ==
"avx512.kandn.w") {
3608 LHS = Builder.CreateNot(
LHS);
3609 Rep = Builder.CreateAnd(
LHS,
RHS);
3610 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3611 }
else if (Name ==
"avx512.kor.w") {
3614 Rep = Builder.CreateOr(
LHS,
RHS);
3615 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3616 }
else if (Name ==
"avx512.kxor.w") {
3619 Rep = Builder.CreateXor(
LHS,
RHS);
3620 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3621 }
else if (Name ==
"avx512.kxnor.w") {
3624 LHS = Builder.CreateNot(
LHS);
3625 Rep = Builder.CreateXor(
LHS,
RHS);
3626 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3627 }
else if (Name ==
"avx512.knot.w") {
3629 Rep = Builder.CreateNot(Rep);
3630 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3631 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3634 Rep = Builder.CreateOr(
LHS,
RHS);
3635 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3637 if (Name[14] ==
'c')
3641 Rep = Builder.CreateICmpEQ(Rep,
C);
3642 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3643 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3644 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3645 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3646 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3649 ConstantInt::get(I32Ty, 0));
3651 ConstantInt::get(I32Ty, 0));
3653 if (Name.contains(
".add."))
3654 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3655 else if (Name.contains(
".sub."))
3656 EltOp = Builder.CreateFSub(Elt0, Elt1);
3657 else if (Name.contains(
".mul."))
3658 EltOp = Builder.CreateFMul(Elt0, Elt1);
3660 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3661 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3662 ConstantInt::get(I32Ty, 0));
3663 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3665 bool CmpEq = Name[16] ==
'e';
3667 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3669 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3676 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3679 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3682 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3689 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3691 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3692 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3694 if (VecWidth == 128 && EltWidth == 32)
3695 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3696 else if (VecWidth == 256 && EltWidth == 32)
3697 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3698 else if (VecWidth == 512 && EltWidth == 32)
3699 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3700 else if (VecWidth == 128 && EltWidth == 64)
3701 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3702 else if (VecWidth == 256 && EltWidth == 64)
3703 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3704 else if (VecWidth == 512 && EltWidth == 64)
3705 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3712 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3715 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3716 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3718 if (VecWidth == 128 && EltWidth == 32)
3719 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3720 else if (VecWidth == 256 && EltWidth == 32)
3721 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3722 else if (VecWidth == 512 && EltWidth == 32)
3723 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3724 else if (VecWidth == 128 && EltWidth == 64)
3725 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3726 else if (VecWidth == 256 && EltWidth == 64)
3727 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3728 else if (VecWidth == 512 && EltWidth == 64)
3729 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3734 if (VecWidth == 512)
3736 Args.push_back(Mask);
3738 Rep = Builder.CreateIntrinsic(IID, Args);
3739 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3743 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3746 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3747 Name.starts_with(
"avx512.cvtw2mask.") ||
3748 Name.starts_with(
"avx512.cvtd2mask.") ||
3749 Name.starts_with(
"avx512.cvtq2mask.")) {
3754 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3755 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3756 Name.starts_with(
"avx512.mask.pabs")) {
3758 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3759 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3760 Name.starts_with(
"avx512.mask.pmaxs")) {
3762 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3763 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3764 Name.starts_with(
"avx512.mask.pmaxu")) {
3766 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3767 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3768 Name.starts_with(
"avx512.mask.pmins")) {
3770 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3771 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3772 Name.starts_with(
"avx512.mask.pminu")) {
3774 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3775 Name ==
"avx512.pmulu.dq.512" ||
3776 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3778 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3779 Name ==
"avx512.pmul.dq.512" ||
3780 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3782 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3783 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3788 }
else if (Name ==
"avx512.cvtusi2sd") {
3793 }
else if (Name ==
"sse2.cvtss2sd") {
3795 Rep = Builder.CreateFPExt(
3798 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3799 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3800 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3801 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3802 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3803 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3804 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3805 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3806 Name ==
"avx512.mask.cvtqq2ps.256" ||
3807 Name ==
"avx512.mask.cvtqq2ps.512" ||
3808 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3809 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3810 Name ==
"avx.cvt.ps2.pd.256" ||
3811 Name ==
"avx512.mask.cvtps2pd.128" ||
3812 Name ==
"avx512.mask.cvtps2pd.256") {
3817 unsigned NumDstElts = DstTy->getNumElements();
3818 if (NumDstElts < SrcTy->getNumElements()) {
3819 assert(NumDstElts == 2 &&
"Unexpected vector size");
3820 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3823 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3824 bool IsUnsigned = Name.contains(
"cvtu");
3826 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3830 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3831 : Intrinsic::x86_avx512_sitofp_round;
3832 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3835 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3836 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3842 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3843 Name.starts_with(
"vcvtph2ps.")) {
3847 unsigned NumDstElts = DstTy->getNumElements();
3848 if (NumDstElts != SrcTy->getNumElements()) {
3849 assert(NumDstElts == 4 &&
"Unexpected vector size");
3850 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3852 Rep = Builder.CreateBitCast(
3854 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3858 }
else if (Name.starts_with(
"avx512.mask.load")) {
3860 bool Aligned = Name[16] !=
'u';
3863 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3867 ResultTy->getNumElements());
3868 Rep = Builder.CreateIntrinsic(
3869 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3871 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3877 Rep = Builder.CreateIntrinsic(
3878 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3880 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3881 Name.starts_with(
"avx512.mask.expand.")) {
3885 ResultTy->getNumElements());
3887 bool IsCompress = Name[12] ==
'c';
3888 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3889 : Intrinsic::x86_avx512_mask_expand;
3890 Rep = Builder.CreateIntrinsic(
3892 }
else if (Name.starts_with(
"xop.vpcom")) {
3894 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3895 Name.ends_with(
"uq"))
3897 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3898 Name.ends_with(
"d") || Name.ends_with(
"q"))
3907 Name = Name.substr(9);
3908 if (Name.starts_with(
"lt"))
3910 else if (Name.starts_with(
"le"))
3912 else if (Name.starts_with(
"gt"))
3914 else if (Name.starts_with(
"ge"))
3916 else if (Name.starts_with(
"eq"))
3918 else if (Name.starts_with(
"ne"))
3920 else if (Name.starts_with(
"false"))
3922 else if (Name.starts_with(
"true"))
3929 }
else if (Name.starts_with(
"xop.vpcmov")) {
3931 Value *NotSel = Builder.CreateNot(Sel);
3934 Rep = Builder.CreateOr(Sel0, Sel1);
3935 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3936 Name.starts_with(
"avx512.mask.prol")) {
3938 }
else if (Name.starts_with(
"avx512.pror") ||
3939 Name.starts_with(
"avx512.mask.pror")) {
3941 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3942 Name.starts_with(
"avx512.mask.vpshld") ||
3943 Name.starts_with(
"avx512.maskz.vpshld")) {
3944 bool ZeroMask = Name[11] ==
'z';
3946 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3947 Name.starts_with(
"avx512.mask.vpshrd") ||
3948 Name.starts_with(
"avx512.maskz.vpshrd")) {
3949 bool ZeroMask = Name[11] ==
'z';
3951 }
else if (Name ==
"sse42.crc32.64.8") {
3954 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3956 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3957 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3958 Name.starts_with(
"avx512.vbroadcast.s")) {
3961 Type *EltTy = VecTy->getElementType();
3962 unsigned EltNum = VecTy->getNumElements();
3966 for (
unsigned I = 0;
I < EltNum; ++
I)
3967 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3968 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3969 Name.starts_with(
"sse41.pmovzx") ||
3970 Name.starts_with(
"avx2.pmovsx") ||
3971 Name.starts_with(
"avx2.pmovzx") ||
3972 Name.starts_with(
"avx512.mask.pmovsx") ||
3973 Name.starts_with(
"avx512.mask.pmovzx")) {
3975 unsigned NumDstElts = DstTy->getNumElements();
3979 for (
unsigned i = 0; i != NumDstElts; ++i)
3984 bool DoSext = Name.contains(
"pmovsx");
3986 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3991 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3992 Name ==
"avx512.mask.pmov.qd.512" ||
3993 Name ==
"avx512.mask.pmov.wb.256" ||
3994 Name ==
"avx512.mask.pmov.wb.512") {
3999 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4000 Name ==
"avx2.vbroadcasti128") {
4006 if (NumSrcElts == 2)
4009 Rep = Builder.CreateShuffleVector(
Load,
4011 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4012 Name.starts_with(
"avx512.mask.shuf.f")) {
4017 unsigned ControlBitsMask = NumLanes - 1;
4018 unsigned NumControlBits = NumLanes / 2;
4021 for (
unsigned l = 0; l != NumLanes; ++l) {
4022 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4024 if (l >= NumLanes / 2)
4025 LaneMask += NumLanes;
4026 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4027 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4033 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4034 Name.starts_with(
"avx512.mask.broadcasti")) {
4037 unsigned NumDstElts =
4041 for (
unsigned i = 0; i != NumDstElts; ++i)
4042 ShuffleMask[i] = i % NumSrcElts;
4048 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4049 Name.starts_with(
"avx2.vbroadcast") ||
4050 Name.starts_with(
"avx512.pbroadcast") ||
4051 Name.starts_with(
"avx512.mask.broadcast.s")) {
4058 Rep = Builder.CreateShuffleVector(
Op, M);
4063 }
else if (Name.starts_with(
"sse2.padds.") ||
4064 Name.starts_with(
"avx2.padds.") ||
4065 Name.starts_with(
"avx512.padds.") ||
4066 Name.starts_with(
"avx512.mask.padds.")) {
4068 }
else if (Name.starts_with(
"sse2.psubs.") ||
4069 Name.starts_with(
"avx2.psubs.") ||
4070 Name.starts_with(
"avx512.psubs.") ||
4071 Name.starts_with(
"avx512.mask.psubs.")) {
4073 }
else if (Name.starts_with(
"sse2.paddus.") ||
4074 Name.starts_with(
"avx2.paddus.") ||
4075 Name.starts_with(
"avx512.mask.paddus.")) {
4077 }
else if (Name.starts_with(
"sse2.psubus.") ||
4078 Name.starts_with(
"avx2.psubus.") ||
4079 Name.starts_with(
"avx512.mask.psubus.")) {
4081 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4086 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4090 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4095 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4100 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4101 Name ==
"avx512.psll.dq.512") {
4105 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4106 Name ==
"avx512.psrl.dq.512") {
4110 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4111 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4112 Name.starts_with(
"avx2.pblendd.")) {
4117 unsigned NumElts = VecTy->getNumElements();
4120 for (
unsigned i = 0; i != NumElts; ++i)
4121 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4123 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4124 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4125 Name ==
"avx2.vinserti128" ||
4126 Name.starts_with(
"avx512.mask.insert")) {
4130 unsigned DstNumElts =
4132 unsigned SrcNumElts =
4134 unsigned Scale = DstNumElts / SrcNumElts;
4141 for (
unsigned i = 0; i != SrcNumElts; ++i)
4143 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4144 Idxs[i] = SrcNumElts;
4145 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4159 for (
unsigned i = 0; i != DstNumElts; ++i)
4162 for (
unsigned i = 0; i != SrcNumElts; ++i)
4163 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4164 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4170 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4171 Name ==
"avx2.vextracti128" ||
4172 Name.starts_with(
"avx512.mask.vextract")) {
4175 unsigned DstNumElts =
4177 unsigned SrcNumElts =
4179 unsigned Scale = SrcNumElts / DstNumElts;
4186 for (
unsigned i = 0; i != DstNumElts; ++i) {
4187 Idxs[i] = i + (
Imm * DstNumElts);
4189 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4195 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4196 Name.starts_with(
"avx512.mask.perm.di.")) {
4200 unsigned NumElts = VecTy->getNumElements();
4203 for (
unsigned i = 0; i != NumElts; ++i)
4204 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4206 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4211 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4223 unsigned HalfSize = NumElts / 2;
4235 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4236 for (
unsigned i = 0; i < HalfSize; ++i)
4237 ShuffleMask[i] = StartIndex + i;
4240 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4241 for (
unsigned i = 0; i < HalfSize; ++i)
4242 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4244 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4246 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4247 Name.starts_with(
"avx512.mask.vpermil.p") ||
4248 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4252 unsigned NumElts = VecTy->getNumElements();
4254 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4255 unsigned IdxMask = ((1 << IdxSize) - 1);
4261 for (
unsigned i = 0; i != NumElts; ++i)
4262 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4264 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4269 }
else if (Name ==
"sse2.pshufl.w" ||
4270 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4275 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4279 for (
unsigned l = 0; l != NumElts; l += 8) {
4280 for (
unsigned i = 0; i != 4; ++i)
4281 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4282 for (
unsigned i = 4; i != 8; ++i)
4283 Idxs[i + l] = i + l;
4286 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4291 }
else if (Name ==
"sse2.pshufh.w" ||
4292 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4297 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4301 for (
unsigned l = 0; l != NumElts; l += 8) {
4302 for (
unsigned i = 0; i != 4; ++i)
4303 Idxs[i + l] = i + l;
4304 for (
unsigned i = 0; i != 4; ++i)
4305 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4308 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4313 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4320 unsigned HalfLaneElts = NumLaneElts / 2;
4323 for (
unsigned i = 0; i != NumElts; ++i) {
4325 Idxs[i] = i - (i % NumLaneElts);
4327 if ((i % NumLaneElts) >= HalfLaneElts)
4331 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4334 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4338 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4339 Name.starts_with(
"avx512.mask.movshdup") ||
4340 Name.starts_with(
"avx512.mask.movsldup")) {
4346 if (Name.starts_with(
"avx512.mask.movshdup."))
4350 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4351 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4352 Idxs[i + l + 0] = i + l +
Offset;
4353 Idxs[i + l + 1] = i + l +
Offset;
4356 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4360 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4361 Name.starts_with(
"avx512.mask.unpckl.")) {
4368 for (
int l = 0; l != NumElts; l += NumLaneElts)
4369 for (
int i = 0; i != NumLaneElts; ++i)
4370 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4372 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4376 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4377 Name.starts_with(
"avx512.mask.unpckh.")) {
4384 for (
int l = 0; l != NumElts; l += NumLaneElts)
4385 for (
int i = 0; i != NumLaneElts; ++i)
4386 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4388 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4392 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4393 Name.starts_with(
"avx512.mask.pand.")) {
4396 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4398 Rep = Builder.CreateBitCast(Rep, FTy);
4401 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4402 Name.starts_with(
"avx512.mask.pandn.")) {
4405 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4406 Rep = Builder.CreateAnd(Rep,
4408 Rep = Builder.CreateBitCast(Rep, FTy);
4411 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4412 Name.starts_with(
"avx512.mask.por.")) {
4415 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4417 Rep = Builder.CreateBitCast(Rep, FTy);
4420 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4421 Name.starts_with(
"avx512.mask.pxor.")) {
4424 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4426 Rep = Builder.CreateBitCast(Rep, FTy);
4429 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4433 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4437 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4441 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4442 if (Name.ends_with(
".512")) {
4444 if (Name[17] ==
's')
4445 IID = Intrinsic::x86_avx512_add_ps_512;
4447 IID = Intrinsic::x86_avx512_add_pd_512;
4449 Rep = Builder.CreateIntrinsic(
4457 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4458 if (Name.ends_with(
".512")) {
4460 if (Name[17] ==
's')
4461 IID = Intrinsic::x86_avx512_div_ps_512;
4463 IID = Intrinsic::x86_avx512_div_pd_512;
4465 Rep = Builder.CreateIntrinsic(
4473 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4474 if (Name.ends_with(
".512")) {
4476 if (Name[17] ==
's')
4477 IID = Intrinsic::x86_avx512_mul_ps_512;
4479 IID = Intrinsic::x86_avx512_mul_pd_512;
4481 Rep = Builder.CreateIntrinsic(
4489 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4490 if (Name.ends_with(
".512")) {
4492 if (Name[17] ==
's')
4493 IID = Intrinsic::x86_avx512_sub_ps_512;
4495 IID = Intrinsic::x86_avx512_sub_pd_512;
4497 Rep = Builder.CreateIntrinsic(
4505 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4506 Name.starts_with(
"avx512.mask.min.p")) &&
4507 Name.drop_front(18) ==
".512") {
4508 bool IsDouble = Name[17] ==
'd';
4509 bool IsMin = Name[13] ==
'i';
4511 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4512 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4515 Rep = Builder.CreateIntrinsic(
4520 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4522 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4523 {CI->getArgOperand(0), Builder.getInt1(false)});
4526 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4527 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4528 bool IsVariable = Name[16] ==
'v';
4529 char Size = Name[16] ==
'.' ? Name[17]
4530 : Name[17] ==
'.' ? Name[18]
4531 : Name[18] ==
'.' ? Name[19]
4535 if (IsVariable && Name[17] !=
'.') {
4536 if (
Size ==
'd' && Name[17] ==
'2')
4537 IID = Intrinsic::x86_avx2_psllv_q;
4538 else if (
Size ==
'd' && Name[17] ==
'4')
4539 IID = Intrinsic::x86_avx2_psllv_q_256;
4540 else if (
Size ==
's' && Name[17] ==
'4')
4541 IID = Intrinsic::x86_avx2_psllv_d;
4542 else if (
Size ==
's' && Name[17] ==
'8')
4543 IID = Intrinsic::x86_avx2_psllv_d_256;
4544 else if (
Size ==
'h' && Name[17] ==
'8')
4545 IID = Intrinsic::x86_avx512_psllv_w_128;
4546 else if (
Size ==
'h' && Name[17] ==
'1')
4547 IID = Intrinsic::x86_avx512_psllv_w_256;
4548 else if (Name[17] ==
'3' && Name[18] ==
'2')
4549 IID = Intrinsic::x86_avx512_psllv_w_512;
4552 }
else if (Name.ends_with(
".128")) {
4554 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4555 : Intrinsic::x86_sse2_psll_d;
4556 else if (
Size ==
'q')
4557 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4558 : Intrinsic::x86_sse2_psll_q;
4559 else if (
Size ==
'w')
4560 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4561 : Intrinsic::x86_sse2_psll_w;
4564 }
else if (Name.ends_with(
".256")) {
4566 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4567 : Intrinsic::x86_avx2_psll_d;
4568 else if (
Size ==
'q')
4569 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4570 : Intrinsic::x86_avx2_psll_q;
4571 else if (
Size ==
'w')
4572 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4573 : Intrinsic::x86_avx2_psll_w;
4578 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4579 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4580 : Intrinsic::x86_avx512_psll_d_512;
4581 else if (
Size ==
'q')
4582 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4583 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4584 : Intrinsic::x86_avx512_psll_q_512;
4585 else if (
Size ==
'w')
4586 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4587 : Intrinsic::x86_avx512_psll_w_512;
4593 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4594 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4595 bool IsVariable = Name[16] ==
'v';
4596 char Size = Name[16] ==
'.' ? Name[17]
4597 : Name[17] ==
'.' ? Name[18]
4598 : Name[18] ==
'.' ? Name[19]
4602 if (IsVariable && Name[17] !=
'.') {
4603 if (
Size ==
'd' && Name[17] ==
'2')
4604 IID = Intrinsic::x86_avx2_psrlv_q;
4605 else if (
Size ==
'd' && Name[17] ==
'4')
4606 IID = Intrinsic::x86_avx2_psrlv_q_256;
4607 else if (
Size ==
's' && Name[17] ==
'4')
4608 IID = Intrinsic::x86_avx2_psrlv_d;
4609 else if (
Size ==
's' && Name[17] ==
'8')
4610 IID = Intrinsic::x86_avx2_psrlv_d_256;
4611 else if (
Size ==
'h' && Name[17] ==
'8')
4612 IID = Intrinsic::x86_avx512_psrlv_w_128;
4613 else if (
Size ==
'h' && Name[17] ==
'1')
4614 IID = Intrinsic::x86_avx512_psrlv_w_256;
4615 else if (Name[17] ==
'3' && Name[18] ==
'2')
4616 IID = Intrinsic::x86_avx512_psrlv_w_512;
4619 }
else if (Name.ends_with(
".128")) {
4621 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4622 : Intrinsic::x86_sse2_psrl_d;
4623 else if (
Size ==
'q')
4624 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4625 : Intrinsic::x86_sse2_psrl_q;
4626 else if (
Size ==
'w')
4627 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4628 : Intrinsic::x86_sse2_psrl_w;
4631 }
else if (Name.ends_with(
".256")) {
4633 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4634 : Intrinsic::x86_avx2_psrl_d;
4635 else if (
Size ==
'q')
4636 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4637 : Intrinsic::x86_avx2_psrl_q;
4638 else if (
Size ==
'w')
4639 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4640 : Intrinsic::x86_avx2_psrl_w;
4645 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4646 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4647 : Intrinsic::x86_avx512_psrl_d_512;
4648 else if (
Size ==
'q')
4649 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4650 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4651 : Intrinsic::x86_avx512_psrl_q_512;
4652 else if (
Size ==
'w')
4653 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4654 : Intrinsic::x86_avx512_psrl_w_512;
4660 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4661 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4662 bool IsVariable = Name[16] ==
'v';
4663 char Size = Name[16] ==
'.' ? Name[17]
4664 : Name[17] ==
'.' ? Name[18]
4665 : Name[18] ==
'.' ? Name[19]
4669 if (IsVariable && Name[17] !=
'.') {
4670 if (
Size ==
's' && Name[17] ==
'4')
4671 IID = Intrinsic::x86_avx2_psrav_d;
4672 else if (
Size ==
's' && Name[17] ==
'8')
4673 IID = Intrinsic::x86_avx2_psrav_d_256;
4674 else if (
Size ==
'h' && Name[17] ==
'8')
4675 IID = Intrinsic::x86_avx512_psrav_w_128;
4676 else if (
Size ==
'h' && Name[17] ==
'1')
4677 IID = Intrinsic::x86_avx512_psrav_w_256;
4678 else if (Name[17] ==
'3' && Name[18] ==
'2')
4679 IID = Intrinsic::x86_avx512_psrav_w_512;
4682 }
else if (Name.ends_with(
".128")) {
4684 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4685 : Intrinsic::x86_sse2_psra_d;
4686 else if (
Size ==
'q')
4687 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4688 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4689 : Intrinsic::x86_avx512_psra_q_128;
4690 else if (
Size ==
'w')
4691 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4692 : Intrinsic::x86_sse2_psra_w;
4695 }
else if (Name.ends_with(
".256")) {
4697 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4698 : Intrinsic::x86_avx2_psra_d;
4699 else if (
Size ==
'q')
4700 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4701 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4702 : Intrinsic::x86_avx512_psra_q_256;
4703 else if (
Size ==
'w')
4704 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4705 : Intrinsic::x86_avx2_psra_w;
4710 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4711 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4712 : Intrinsic::x86_avx512_psra_d_512;
4713 else if (
Size ==
'q')
4714 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4715 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4716 : Intrinsic::x86_avx512_psra_q_512;
4717 else if (
Size ==
'w')
4718 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4719 : Intrinsic::x86_avx512_psra_w_512;
4725 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4727 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4729 }
else if (Name.ends_with(
".movntdqa")) {
4733 LoadInst *LI = Builder.CreateAlignedLoad(
4738 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4739 Name.starts_with(
"fma.vfmsub.") ||
4740 Name.starts_with(
"fma.vfnmadd.") ||
4741 Name.starts_with(
"fma.vfnmsub.")) {
4742 bool NegMul = Name[6] ==
'n';
4743 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4744 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4755 if (NegMul && !IsScalar)
4756 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4757 if (NegMul && IsScalar)
4758 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4760 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4762 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4766 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4774 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4778 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4779 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4780 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4781 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4782 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4783 bool IsMask3 = Name[11] ==
'3';
4784 bool IsMaskZ = Name[11] ==
'z';
4786 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4787 bool NegMul = Name[2] ==
'n';
4788 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4794 if (NegMul && (IsMask3 || IsMaskZ))
4795 A = Builder.CreateFNeg(
A);
4796 if (NegMul && !(IsMask3 || IsMaskZ))
4797 B = Builder.CreateFNeg(
B);
4799 C = Builder.CreateFNeg(
C);
4801 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4802 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4803 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4810 if (Name.back() ==
'd')
4811 IID = Intrinsic::x86_avx512_vfmadd_f64;
4813 IID = Intrinsic::x86_avx512_vfmadd_f32;
4814 Rep = Builder.CreateIntrinsic(IID,
Ops);
4816 Rep = Builder.CreateFMA(
A,
B,
C);
4825 if (NegAcc && IsMask3)
4830 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4832 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4833 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4834 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4835 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4836 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4837 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4838 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4839 bool IsMask3 = Name[11] ==
'3';
4840 bool IsMaskZ = Name[11] ==
'z';
4842 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4843 bool NegMul = Name[2] ==
'n';
4844 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4850 if (NegMul && (IsMask3 || IsMaskZ))
4851 A = Builder.CreateFNeg(
A);
4852 if (NegMul && !(IsMask3 || IsMaskZ))
4853 B = Builder.CreateFNeg(
B);
4855 C = Builder.CreateFNeg(
C);
4862 if (Name[Name.size() - 5] ==
's')
4863 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4865 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4869 Rep = Builder.CreateFMA(
A,
B,
C);
4877 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4881 if (VecWidth == 128 && EltWidth == 32)
4882 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4883 else if (VecWidth == 256 && EltWidth == 32)
4884 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4885 else if (VecWidth == 128 && EltWidth == 64)
4886 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4887 else if (VecWidth == 256 && EltWidth == 64)
4888 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4894 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4895 Rep = Builder.CreateIntrinsic(IID,
Ops);
4896 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4897 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4898 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4899 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4900 bool IsMask3 = Name[11] ==
'3';
4901 bool IsMaskZ = Name[11] ==
'z';
4903 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4904 bool IsSubAdd = Name[3] ==
's';
4908 if (Name[Name.size() - 5] ==
's')
4909 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4911 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4916 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4918 Rep = Builder.CreateIntrinsic(IID,
Ops);
4927 Value *Odd = Builder.CreateCall(FMA,
Ops);
4928 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4929 Value *Even = Builder.CreateCall(FMA,
Ops);
4935 for (
int i = 0; i != NumElts; ++i)
4936 Idxs[i] = i + (i % 2) * NumElts;
4938 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4946 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4947 Name.starts_with(
"avx512.maskz.pternlog.")) {
4948 bool ZeroMask = Name[11] ==
'z';
4952 if (VecWidth == 128 && EltWidth == 32)
4953 IID = Intrinsic::x86_avx512_pternlog_d_128;
4954 else if (VecWidth == 256 && EltWidth == 32)
4955 IID = Intrinsic::x86_avx512_pternlog_d_256;
4956 else if (VecWidth == 512 && EltWidth == 32)
4957 IID = Intrinsic::x86_avx512_pternlog_d_512;
4958 else if (VecWidth == 128 && EltWidth == 64)
4959 IID = Intrinsic::x86_avx512_pternlog_q_128;
4960 else if (VecWidth == 256 && EltWidth == 64)
4961 IID = Intrinsic::x86_avx512_pternlog_q_256;
4962 else if (VecWidth == 512 && EltWidth == 64)
4963 IID = Intrinsic::x86_avx512_pternlog_q_512;
4969 Rep = Builder.CreateIntrinsic(IID, Args);
4973 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4974 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4975 bool ZeroMask = Name[11] ==
'z';
4976 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4979 if (VecWidth == 128 && !
High)
4980 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4981 else if (VecWidth == 256 && !
High)
4982 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4983 else if (VecWidth == 512 && !
High)
4984 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4985 else if (VecWidth == 128 &&
High)
4986 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4987 else if (VecWidth == 256 &&
High)
4988 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4989 else if (VecWidth == 512 &&
High)
4990 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4996 Rep = Builder.CreateIntrinsic(IID, Args);
5000 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5001 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5002 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5003 bool ZeroMask = Name[11] ==
'z';
5004 bool IndexForm = Name[17] ==
'i';
5006 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5007 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5008 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5009 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5010 bool ZeroMask = Name[11] ==
'z';
5011 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5014 if (VecWidth == 128 && !IsSaturating)
5015 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5016 else if (VecWidth == 256 && !IsSaturating)
5017 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5018 else if (VecWidth == 512 && !IsSaturating)
5019 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5020 else if (VecWidth == 128 && IsSaturating)
5021 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5022 else if (VecWidth == 256 && IsSaturating)
5023 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5024 else if (VecWidth == 512 && IsSaturating)
5025 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5035 if (Args[1]->
getType()->isVectorTy() &&
5038 ->isIntegerTy(32) &&
5039 Args[2]->
getType()->isVectorTy() &&
5042 ->isIntegerTy(32)) {
5043 Type *NewArgType =
nullptr;
5044 if (VecWidth == 128)
5046 else if (VecWidth == 256)
5048 else if (VecWidth == 512)
5054 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5055 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5058 Rep = Builder.CreateIntrinsic(IID, Args);
5062 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5063 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5064 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5065 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5066 bool ZeroMask = Name[11] ==
'z';
5067 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5070 if (VecWidth == 128 && !IsSaturating)
5071 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5072 else if (VecWidth == 256 && !IsSaturating)
5073 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5074 else if (VecWidth == 512 && !IsSaturating)
5075 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5076 else if (VecWidth == 128 && IsSaturating)
5077 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5078 else if (VecWidth == 256 && IsSaturating)
5079 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5080 else if (VecWidth == 512 && IsSaturating)
5081 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5091 if (Args[1]->
getType()->isVectorTy() &&
5094 ->isIntegerTy(32) &&
5095 Args[2]->
getType()->isVectorTy() &&
5098 ->isIntegerTy(32)) {
5099 Type *NewArgType =
nullptr;
5100 if (VecWidth == 128)
5102 else if (VecWidth == 256)
5104 else if (VecWidth == 512)
5110 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5111 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5114 Rep = Builder.CreateIntrinsic(IID, Args);
5118 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5119 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5120 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5122 if (Name[0] ==
'a' && Name.back() ==
'2')
5123 IID = Intrinsic::x86_addcarry_32;
5124 else if (Name[0] ==
'a' && Name.back() ==
'4')
5125 IID = Intrinsic::x86_addcarry_64;
5126 else if (Name[0] ==
's' && Name.back() ==
'2')
5127 IID = Intrinsic::x86_subborrow_32;
5128 else if (Name[0] ==
's' && Name.back() ==
'4')
5129 IID = Intrinsic::x86_subborrow_64;
5136 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5139 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5142 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5146 }
else if (Name.starts_with(
"avx512.mask.") &&
5149 }
else if (Name.starts_with(
"bmi.pdep.")) {
5151 }
else if (Name.starts_with(
"bmi.pext.")) {
5161 if (Name.starts_with(
"neon.bfcvt")) {
5162 if (Name.starts_with(
"neon.bfcvtn2")) {
5164 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5166 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5167 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5170 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5171 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5173 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5177 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5178 return Builder.CreateShuffleVector(
5181 return Builder.CreateFPTrunc(CI->
getOperand(0),
5184 }
else if (Name.starts_with(
"sve.fcvt")) {
5187 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5188 .
Case(
"sve.fcvtnt.bf16f32",
5189 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5201 if (Args[1]->
getType() != BadPredTy)
5204 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5205 BadPredTy, Args[1]);
5206 Args[1] = Builder.CreateIntrinsic(
5207 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5209 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5213 if (Name ==
"neon.vcvtfp2hf")
5214 return Builder.CreateBitCast(
5215 Builder.CreateFPTrunc(
5219 if (Name ==
"neon.vcvthf2fp")
5220 return Builder.CreateFPExt(
5221 Builder.CreateBitCast(
5231 if (Name ==
"mve.vctp64.old") {
5234 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5237 Value *C1 = Builder.CreateIntrinsic(
5238 Intrinsic::arm_mve_pred_v2i,
5240 return Builder.CreateIntrinsic(
5241 Intrinsic::arm_mve_pred_i2v,
5243 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5244 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5245 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5246 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5248 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5249 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5250 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5251 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5253 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5254 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5255 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5256 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5257 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5258 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5259 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5260 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5261 std::vector<Type *> Tys;
5265 case Intrinsic::arm_mve_mull_int_predicated:
5266 case Intrinsic::arm_mve_vqdmull_predicated:
5267 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5270 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5271 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5272 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5276 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5280 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5284 case Intrinsic::arm_cde_vcx1q_predicated:
5285 case Intrinsic::arm_cde_vcx1qa_predicated:
5286 case Intrinsic::arm_cde_vcx2q_predicated:
5287 case Intrinsic::arm_cde_vcx2qa_predicated:
5288 case Intrinsic::arm_cde_vcx3q_predicated:
5289 case Intrinsic::arm_cde_vcx3qa_predicated:
5296 std::vector<Value *>
Ops;
5298 Type *Ty =
Op->getType();
5299 if (Ty->getScalarSizeInBits() == 1) {
5300 Value *C1 = Builder.CreateIntrinsic(
5301 Intrinsic::arm_mve_pred_v2i,
5303 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5308 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5323 auto UpgradeLegacyWMMAIUIntrinsicCall =
5328 Args.push_back(Builder.getFalse());
5332 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5339 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5343 NewCall->copyMetadata(*CI);
5347 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5348 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5349 "intrinsic should have 7 arguments");
5352 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5354 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5355 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5356 "intrinsic should have 8 arguments");
5361 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5364 switch (
F->getIntrinsicID()) {
5367 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5368 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5369 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5370 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5371 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5372 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5387 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5390 F->getParent(),
F->getIntrinsicID(), Overloads);
5395 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5399 NewCall->copyMetadata(*CI);
5400 NewCall->takeName(CI);
5405 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5411 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5412 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5420 if (Name.starts_with(
"addrspacecast.nonnull")) {
5423 Value *ASC = Builder.CreateAddrSpaceCast(
5446 if (NumOperands < 3)
5459 bool IsVolatile =
false;
5463 if (NumOperands > 3)
5468 if (NumOperands > 5) {
5470 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5484 if (VT->getElementType()->isIntegerTy(16)) {
5487 Val = Builder.CreateBitCast(Val, AsBF16);
5495 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5497 unsigned AddrSpace = PtrTy->getAddressSpace();
5500 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5502 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5507 MDNode *RangeNotPrivate =
5510 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5516 return Builder.CreateBitCast(RMW, RetTy);
5537 return MAV->getMetadata();
5546 if (Name ==
"label") {
5548 }
else if (Name ==
"assign") {
5555 }
else if (Name ==
"declare") {
5559 }
else if (Name ==
"addr") {
5569 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5570 }
else if (Name ==
"value") {
5573 unsigned ExprOp = 2;
5588 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5596 int64_t OffsetVal =
Offset->getSExtValue();
5597 return Builder.CreateIntrinsic(OffsetVal >= 0
5598 ? Intrinsic::vector_splice_left
5599 : Intrinsic::vector_splice_right,
5601 {CI->getArgOperand(0), CI->getArgOperand(1),
5602 Builder.getInt32(std::abs(OffsetVal))});
5607 if (Name.starts_with(
"to.fp16")) {
5609 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5610 return Builder.CreateBitCast(Cast, CI->
getType());
5613 if (Name.starts_with(
"from.fp16")) {
5615 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5616 return Builder.CreateFPExt(Cast, CI->
getType());
5675 else if (Opcode == Instruction::ICmp)
5678 else if (Opcode == Instruction::FCmp)
5681 else if (Opcode == Instruction::Select)
5686 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5698 if (Defaults.empty())
5701 unsigned OldArgCount = CI->
arg_size();
5702 unsigned NewArgCount = NewFn->
arg_size();
5704 if (OldArgCount < FirstDefault)
5708 if (OldArgCount > NewArgCount)
5713 if (OldArgCount == NewArgCount) {
5725 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5726 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5727 "missing argument outside the default range");
5728 Type *ParamTy = NewFT->getParamType(Idx);
5733 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5739 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5771 if (!Name.consume_front(
"llvm."))
5774 bool IsX86 = Name.consume_front(
"x86.");
5775 bool IsNVVM = Name.consume_front(
"nvvm.");
5776 bool IsAArch64 = Name.consume_front(
"aarch64.");
5777 bool IsARM = Name.consume_front(
"arm.");
5778 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5779 bool IsDbg = Name.consume_front(
"dbg.");
5781 (Name.consume_front(
"experimental.vector.splice") ||
5782 Name.consume_front(
"vector.splice")) &&
5783 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5784 Value *Rep =
nullptr;
5786 if (!IsX86 && Name ==
"stackprotectorcheck") {
5788 }
else if (IsNVVM) {
5792 }
else if (IsAArch64) {
5796 }
else if (IsAMDGCN) {
5800 }
else if (IsOldSplice) {
5802 }
else if (Name.consume_front(
"convert.")) {
5804 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5819 const auto &DefaultCase = [&]() ->
void {
5827 "Unknown function for CallBase upgrade and isn't just a name change");
5835 "Return type must have changed");
5836 assert(OldST->getNumElements() ==
5838 "Must have same number of elements");
5841 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5844 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5845 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5846 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5867 case Intrinsic::arm_neon_vst1:
5868 case Intrinsic::arm_neon_vst2:
5869 case Intrinsic::arm_neon_vst3:
5870 case Intrinsic::arm_neon_vst4:
5871 case Intrinsic::arm_neon_vst2lane:
5872 case Intrinsic::arm_neon_vst3lane:
5873 case Intrinsic::arm_neon_vst4lane: {
5875 NewCall = Builder.CreateCall(NewFn, Args);
5878 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5879 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5880 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5885 NewCall = Builder.CreateCall(NewFn, Args);
5888 case Intrinsic::aarch64_sve_ld3_sret:
5889 case Intrinsic::aarch64_sve_ld4_sret:
5890 case Intrinsic::aarch64_sve_ld2_sret: {
5898 Name = Name.substr(5);
5905 unsigned MinElts = RetTy->getMinNumElements() /
N;
5907 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5909 for (
unsigned I = 0;
I <
N;
I++) {
5910 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5911 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5917 case Intrinsic::coro_end_async:
5918 case Intrinsic::coro_end: {
5920 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5922 NewCall = Builder.CreateCall(NewFn, Args);
5927 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5928 Value *InRamp = Builder.CreateCall(IsInRamp);
5938 case Intrinsic::vector_extract: {
5940 Name = Name.substr(5);
5941 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5946 unsigned MinElts = RetTy->getMinNumElements();
5949 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5953 case Intrinsic::vector_insert: {
5955 Name = Name.substr(5);
5956 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5960 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5965 NewCall = Builder.CreateCall(
5969 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5975 assert(
N > 1 &&
"Create is expected to be between 2-4");
5978 unsigned MinElts = RetTy->getMinNumElements() /
N;
5979 for (
unsigned I = 0;
I <
N;
I++) {
5981 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5988 case Intrinsic::arm_neon_bfdot:
5989 case Intrinsic::arm_neon_bfmmla:
5990 case Intrinsic::arm_neon_bfmlalb:
5991 case Intrinsic::arm_neon_bfmlalt:
5992 case Intrinsic::aarch64_neon_bfdot:
5993 case Intrinsic::aarch64_neon_bfmmla:
5994 case Intrinsic::aarch64_neon_bfmlalb:
5995 case Intrinsic::aarch64_neon_bfmlalt: {
5998 "Mismatch between function args and call args");
5999 size_t OperandWidth =
6001 assert((OperandWidth == 64 || OperandWidth == 128) &&
6002 "Unexpected operand width");
6004 auto Iter = CI->
args().begin();
6005 Args.push_back(*Iter++);
6006 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6007 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6008 NewCall = Builder.CreateCall(NewFn, Args);
6012 case Intrinsic::bitreverse:
6013 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6016 case Intrinsic::ctlz:
6017 case Intrinsic::cttz: {
6024 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6028 case Intrinsic::objectsize: {
6029 Value *NullIsUnknownSize =
6033 NewCall = Builder.CreateCall(
6038 case Intrinsic::ctpop:
6039 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6041 case Intrinsic::dbg_value: {
6043 Name = Name.substr(5);
6045 if (Name.starts_with(
"dbg.addr")) {
6059 if (
Offset->isNullValue()) {
6060 NewCall = Builder.CreateCall(
6069 case Intrinsic::ptr_annotation:
6077 NewCall = Builder.CreateCall(
6086 case Intrinsic::var_annotation:
6093 NewCall = Builder.CreateCall(
6102 case Intrinsic::riscv_aes32dsi:
6103 case Intrinsic::riscv_aes32dsmi:
6104 case Intrinsic::riscv_aes32esi:
6105 case Intrinsic::riscv_aes32esmi:
6106 case Intrinsic::riscv_sm4ks:
6107 case Intrinsic::riscv_sm4ed: {
6117 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6118 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6124 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6125 Value *Res = NewCall;
6127 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6133 case Intrinsic::nvvm_mapa_shared_cluster: {
6137 Value *Res = NewCall;
6138 Res = Builder.CreateAddrSpaceCast(
6145 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6147 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6149 Args[0] = Builder.CreateAddrSpaceCast(
6153 Args.push_back(Builder.getInt32(0));
6155 NewCall = Builder.CreateCall(NewFn, Args);
6161 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6166 for (
unsigned I = 0;
I < 4; ++
I)
6168 Args.push_back(Builder.getInt32(0));
6169 Args.push_back(Builder.getInt32(0));
6172 Args.push_back(Builder.getInt1(
false));
6173 Args.push_back(Builder.getInt32(0));
6175 NewCall = Builder.CreateCall(NewFn, Args);
6181 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6184 Args[0] = Builder.CreateAddrSpaceCast(
6187 NewCall = Builder.CreateCall(NewFn, Args);
6194#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6195 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6197#undef G2S_CLUSTER_CASE
6202 Args[0] = Builder.CreateAddrSpaceCast(
6208 Args.push_back(Builder.getInt32(0));
6210 NewCall = Builder.CreateCall(NewFn, Args);
6217#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6218 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6226 "expected only the trailing flag_valid_pattern to be missing");
6227 Args.push_back(Builder.getInt32(0));
6229 NewCall = Builder.CreateCall(NewFn, Args);
6235#undef NVVM_TMA_G2S_MODES
6238 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6239 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6240 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6241 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6242 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6243 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6244 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6245 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6247 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6251 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6252 NewCall = Builder.CreateCall(NewFn, Args);
6255 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6256 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6257 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6258 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6261 Builder.getFalse()});
6263 case Intrinsic::nvvm_mbarrier_init: {
6267 if (Args.size() == 2)
6268 Args.push_back(Builder.getInt32(0));
6269 NewCall = Builder.CreateCall(NewFn, Args);
6272 case Intrinsic::riscv_sha256sig0:
6273 case Intrinsic::riscv_sha256sig1:
6274 case Intrinsic::riscv_sha256sum0:
6275 case Intrinsic::riscv_sha256sum1:
6276 case Intrinsic::riscv_sm3p0:
6277 case Intrinsic::riscv_sm3p1: {
6284 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6286 NewCall = Builder.CreateCall(NewFn, Arg);
6288 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6295 case Intrinsic::x86_xop_vfrcz_ss:
6296 case Intrinsic::x86_xop_vfrcz_sd:
6297 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6300 case Intrinsic::x86_xop_vpermil2pd:
6301 case Intrinsic::x86_xop_vpermil2ps:
6302 case Intrinsic::x86_xop_vpermil2pd_256:
6303 case Intrinsic::x86_xop_vpermil2ps_256: {
6307 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6308 NewCall = Builder.CreateCall(NewFn, Args);
6312 case Intrinsic::x86_sse41_ptestc:
6313 case Intrinsic::x86_sse41_ptestz:
6314 case Intrinsic::x86_sse41_ptestnzc: {
6328 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6329 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6331 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6335 case Intrinsic::x86_rdtscp: {
6341 NewCall = Builder.CreateCall(NewFn);
6343 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6346 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6354 case Intrinsic::x86_sse41_insertps:
6355 case Intrinsic::x86_sse41_dppd:
6356 case Intrinsic::x86_sse41_dpps:
6357 case Intrinsic::x86_sse41_mpsadbw:
6358 case Intrinsic::x86_avx_dp_ps_256:
6359 case Intrinsic::x86_avx2_mpsadbw: {
6365 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6366 NewCall = Builder.CreateCall(NewFn, Args);
6370 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6371 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6372 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6373 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6374 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6375 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6381 NewCall = Builder.CreateCall(NewFn, Args);
6390 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6391 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6392 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6393 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6394 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6395 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6399 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6400 Args[1] = Builder.CreateBitCast(
6403 NewCall = Builder.CreateCall(NewFn, Args);
6404 Value *Res = Builder.CreateBitCast(
6412 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6413 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6414 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6418 Args[1] = Builder.CreateBitCast(
6420 Args[2] = Builder.CreateBitCast(
6423 NewCall = Builder.CreateCall(NewFn, Args);
6427 case Intrinsic::thread_pointer: {
6428 NewCall = Builder.CreateCall(NewFn, {});
6432 case Intrinsic::memcpy:
6433 case Intrinsic::memmove:
6434 case Intrinsic::memset: {
6450 NewCall = Builder.CreateCall(NewFn, Args);
6453 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6454 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6455 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6460 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6463 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6467 case Intrinsic::masked_load:
6468 case Intrinsic::masked_gather:
6469 case Intrinsic::masked_store:
6470 case Intrinsic::masked_scatter: {
6476 auto GetMaybeAlign = [](
Value *
Op) {
6478 uint64_t Val = CI->getZExtValue();
6486 auto GetAlign = [&](
Value *
Op) {
6495 case Intrinsic::masked_load:
6496 NewCall = Builder.CreateMaskedLoad(
6500 case Intrinsic::masked_gather:
6501 NewCall = Builder.CreateMaskedGather(
6507 case Intrinsic::masked_store:
6508 NewCall = Builder.CreateMaskedStore(
6512 case Intrinsic::masked_scatter:
6513 NewCall = Builder.CreateMaskedScatter(
6515 DL.getValueOrABITypeAlignment(
6529 case Intrinsic::lifetime_start:
6530 case Intrinsic::lifetime_end: {
6542 NewCall = Builder.CreateLifetimeStart(Ptr);
6544 NewCall = Builder.CreateLifetimeEnd(Ptr);
6553 case Intrinsic::x86_avx512_vpdpbusd_128:
6554 case Intrinsic::x86_avx512_vpdpbusd_256:
6555 case Intrinsic::x86_avx512_vpdpbusd_512:
6556 case Intrinsic::x86_avx512_vpdpbusds_128:
6557 case Intrinsic::x86_avx512_vpdpbusds_256:
6558 case Intrinsic::x86_avx512_vpdpbusds_512:
6559 case Intrinsic::x86_avx2_vpdpbssd_128:
6560 case Intrinsic::x86_avx2_vpdpbssd_256:
6561 case Intrinsic::x86_avx10_vpdpbssd_512:
6562 case Intrinsic::x86_avx2_vpdpbssds_128:
6563 case Intrinsic::x86_avx2_vpdpbssds_256:
6564 case Intrinsic::x86_avx10_vpdpbssds_512:
6565 case Intrinsic::x86_avx2_vpdpbsud_128:
6566 case Intrinsic::x86_avx2_vpdpbsud_256:
6567 case Intrinsic::x86_avx10_vpdpbsud_512:
6568 case Intrinsic::x86_avx2_vpdpbsuds_128:
6569 case Intrinsic::x86_avx2_vpdpbsuds_256:
6570 case Intrinsic::x86_avx10_vpdpbsuds_512:
6571 case Intrinsic::x86_avx2_vpdpbuud_128:
6572 case Intrinsic::x86_avx2_vpdpbuud_256:
6573 case Intrinsic::x86_avx10_vpdpbuud_512:
6574 case Intrinsic::x86_avx2_vpdpbuuds_128:
6575 case Intrinsic::x86_avx2_vpdpbuuds_256:
6576 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6581 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6582 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6584 NewCall = Builder.CreateCall(NewFn, Args);
6587 case Intrinsic::x86_avx512_vpdpwssd_128:
6588 case Intrinsic::x86_avx512_vpdpwssd_256:
6589 case Intrinsic::x86_avx512_vpdpwssd_512:
6590 case Intrinsic::x86_avx512_vpdpwssds_128:
6591 case Intrinsic::x86_avx512_vpdpwssds_256:
6592 case Intrinsic::x86_avx512_vpdpwssds_512:
6593 case Intrinsic::x86_avx2_vpdpwsud_128:
6594 case Intrinsic::x86_avx2_vpdpwsud_256:
6595 case Intrinsic::x86_avx10_vpdpwsud_512:
6596 case Intrinsic::x86_avx2_vpdpwsuds_128:
6597 case Intrinsic::x86_avx2_vpdpwsuds_256:
6598 case Intrinsic::x86_avx10_vpdpwsuds_512:
6599 case Intrinsic::x86_avx2_vpdpwusd_128:
6600 case Intrinsic::x86_avx2_vpdpwusd_256:
6601 case Intrinsic::x86_avx10_vpdpwusd_512:
6602 case Intrinsic::x86_avx2_vpdpwusds_128:
6603 case Intrinsic::x86_avx2_vpdpwusds_256:
6604 case Intrinsic::x86_avx10_vpdpwusds_512:
6605 case Intrinsic::x86_avx2_vpdpwuud_128:
6606 case Intrinsic::x86_avx2_vpdpwuud_256:
6607 case Intrinsic::x86_avx10_vpdpwuud_512:
6608 case Intrinsic::x86_avx2_vpdpwuuds_128:
6609 case Intrinsic::x86_avx2_vpdpwuuds_256:
6610 case Intrinsic::x86_avx10_vpdpwuuds_512:
6615 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6616 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6618 NewCall = Builder.CreateCall(NewFn, Args);
6621 assert(NewCall &&
"Should have either set this variable or returned through "
6622 "the default case");
6629 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6643 F->eraseFromParent();
6649 if (NumOperands == 0)
6657 if (NumOperands == 3) {
6661 Metadata *Elts2[] = {ScalarType, ScalarType,
6675 if (
Opc != Instruction::BitCast)
6679 Type *SrcTy = V->getType();
6696 if (
Opc != Instruction::BitCast)
6699 Type *SrcTy =
C->getType();
6716 if (Flag.getNumOperands() < 3)
6717 return std::nullopt;
6719 return Name->getString();
6720 return std::nullopt;
6734 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6735 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6736 if (auto Name = getModuleFlagNameSafely(*Flag))
6737 return *Name ==
"Debug Info Version";
6740 if (OpIt != ModFlags->op_end()) {
6741 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6748 bool BrokenDebugInfo =
false;
6751 if (!BrokenDebugInfo)
6757 M.getContext().diagnose(Diag);
6764 M.getContext().diagnose(DiagVersion);
6774 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6777 if (
F->hasFnAttribute(Attr)) {
6780 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6782 auto [Part, Rest] = S.
split(
',');
6788 const unsigned Dim = DimC -
'x';
6789 assert(Dim < 3 &&
"Unexpected dim char");
6799 F->addFnAttr(Attr, NewAttr);
6803 return S ==
"x" || S ==
"y" || S ==
"z";
6808 if (
K ==
"kernel") {
6820 const unsigned Idx = (AlignIdxValuePair >> 16);
6821 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6826 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6831 if (
K ==
"minctasm") {
6836 if (
K ==
"maxnreg") {
6841 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6845 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6849 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6853 if (
K ==
"grid_constant") {
6868 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6875 if (!SeenNodes.
insert(MD).second)
6882 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6889 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6891 const MDOperand &V = MD->getOperand(j + 1);
6897 if (NewOperands.
size() > 1)
6910 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6911 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6912 if (ModRetainReleaseMarker) {
6918 ID->getString().split(ValueComp,
"#");
6919 if (ValueComp.
size() == 2) {
6920 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6924 M.eraseNamedMetadata(ModRetainReleaseMarker);
6935 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6961 bool InvalidCast =
false;
6963 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6976 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6978 Args.push_back(Arg);
6985 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6990 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7003 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7011 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7012 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7013 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7014 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7015 {
"objc_autoreleaseReturnValue",
7016 llvm::Intrinsic::objc_autoreleaseReturnValue},
7017 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7018 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7019 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7020 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7021 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7022 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7023 {
"objc_release", llvm::Intrinsic::objc_release},
7024 {
"objc_retain", llvm::Intrinsic::objc_retain},
7025 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7026 {
"objc_retainAutoreleaseReturnValue",
7027 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7028 {
"objc_retainAutoreleasedReturnValue",
7029 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7030 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7031 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7032 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7033 {
"objc_unsafeClaimAutoreleasedReturnValue",
7034 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7035 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7036 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7037 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7038 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7039 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7040 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7041 {
"objc_arc_annotation_topdown_bbstart",
7042 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7043 {
"objc_arc_annotation_topdown_bbend",
7044 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7045 {
"objc_arc_annotation_bottomup_bbstart",
7046 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7047 {
"objc_arc_annotation_bottomup_bbend",
7048 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7050 for (
auto &
I : RuntimeFuncs)
7051 UpgradeToIntrinsic(
I.first,
I.second);
7075 std::optional<bool> UseAddressDisc;
7078 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7079 for (
const MDNode *Flag : ModFlags->operands()) {
7081 if (Name && (*Name ==
"ptrauth-init-fini" ||
7082 *Name ==
"ptrauth-init-fini-address-discrimination"))
7087 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7088 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7089 constexpr unsigned ExpectedAddressMarker = 1;
7092 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7095 bool HasAddressDisc;
7096 if (!CPA->hasAddressDiscriminator())
7097 HasAddressDisc =
false;
7098 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7099 HasAddressDisc =
true;
7103 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7106 UseAddressDisc = HasAddressDisc;
7107 return CPA->getPointer();
7111 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7114 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7116 if (!GV || !GV->hasInitializer())
7120 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7123 std::vector<Constant *> NewStructors;
7124 NewStructors.reserve(OldStructorsArray->getNumOperands());
7126 for (
Use &U : OldStructorsArray->operands()) {
7135 Func = UpgradeSinglePointer(Func);
7139 NewStructors.push_back(
7148 if (GlobalArraysToUpgrade.
empty())
7150 assert(UseAddressDisc.has_value());
7152 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7153 GV->setInitializer(NewInit);
7156 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7166 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7170 bool HasObjCFlag =
false, HasClassProperties =
false;
7171 bool HasSwiftVersionFlag =
false;
7172 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7179 if (
Op->getNumOperands() != 3)
7193 if (ID->getString() ==
"Objective-C Image Info Version")
7195 if (ID->getString() ==
"Objective-C Class Properties")
7196 HasClassProperties =
true;
7198 if (ID->getString() ==
"PIC Level") {
7199 if (
auto *Behavior =
7201 uint64_t V = Behavior->getLimitedValue();
7207 if (ID->getString() ==
"PIE Level")
7208 if (
auto *Behavior =
7215 if (ID->getString() ==
"branch-target-enforcement" ||
7216 ID->getString().starts_with(
"sign-return-address")) {
7217 if (
auto *Behavior =
7223 Op->getOperand(1),
Op->getOperand(2)};
7233 if (ID->getString() ==
"Objective-C Image Info Section") {
7236 Value->getString().split(ValueComp,
" ");
7237 if (ValueComp.
size() != 1) {
7238 std::string NewValue;
7239 for (
auto &S : ValueComp)
7240 NewValue += S.str();
7251 if (ID->getString() ==
"Objective-C Garbage Collection") {
7254 assert(Md->getValue() &&
"Expected non-empty metadata");
7255 auto Type = Md->getValue()->getType();
7258 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7259 if ((Val & 0xff) != Val) {
7260 HasSwiftVersionFlag =
true;
7261 SwiftABIVersion = (Val & 0xff00) >> 8;
7262 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7263 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7274 if (ID->getString() ==
"amdgpu_code_object_version") {
7277 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7286 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7315 if (HasObjCFlag && !HasClassProperties) {
7321 if (HasSwiftVersionFlag) {
7325 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7327 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7335 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7339 auto MatchesVersion = [](
const MDNode *
Op) {
7340 return Op->getNumOperands() >= 3 &&
7354 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7355 assert(
Op->getNumOperands() >= 2 &&
7356 "Expected at least 2 operands - name and linkage type");
7368 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7379 auto TrimSpaces = [](
StringRef Section) -> std::string {
7381 Section.split(Components,
',');
7386 for (
auto Component : Components)
7387 OS <<
',' << Component.trim();
7392 for (
auto &GV : M.globals()) {
7393 if (!GV.hasSection())
7398 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7403 GV.setSection(TrimSpaces(Section));
7419struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7420 StrictFPUpgradeVisitor() =
default;
7423 if (!
Call.isStrictFP())
7429 Call.removeFnAttr(Attribute::StrictFP);
7430 Call.addFnAttr(Attribute::NoBuiltin);
7435struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7436 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7437 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7439 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7454 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7455 StrictFPUpgradeVisitor SFPV;
7460 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7461 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7462 for (
auto &Arg :
F.args())
7464 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7466 bool AddingAttrs =
false, RemovingAttrs =
false;
7467 AttrBuilder AttrsToAdd(
F.getContext());
7472 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7473 A.isValid() &&
A.isStringAttribute()) {
7474 F.setSection(
A.getValueAsString());
7476 RemovingAttrs =
true;
7480 A.isValid() &&
A.isStringAttribute()) {
7483 AddingAttrs = RemovingAttrs =
true;
7486 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7487 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7489 RemovingAttrs =
true;
7490 if (
A.getValueAsString() ==
"true") {
7491 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7500 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7503 if (
A.getValueAsBool()) {
7504 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7510 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7511 RemovingAttrs =
true;
7518 bool HandleDenormalMode =
false;
7520 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7523 DenormalFPMath = ParsedMode;
7525 AddingAttrs = RemovingAttrs =
true;
7526 HandleDenormalMode =
true;
7530 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7534 DenormalFPMathF32 = ParsedMode;
7536 AddingAttrs = RemovingAttrs =
true;
7537 HandleDenormalMode =
true;
7541 if (HandleDenormalMode)
7542 AttrsToAdd.addDenormalFPEnvAttr(
7546 F.removeFnAttrs(AttrsToRemove);
7549 F.addFnAttrs(AttrsToAdd);
7555 if (!
F.hasFnAttribute(FnAttrName))
7556 F.addFnAttr(FnAttrName,
Value);
7563 if (!
F.hasFnAttribute(FnAttrName)) {
7565 F.addFnAttr(FnAttrName);
7567 auto A =
F.getFnAttribute(FnAttrName);
7568 if (
"false" ==
A.getValueAsString())
7569 F.removeFnAttr(FnAttrName);
7570 else if (
"true" ==
A.getValueAsString()) {
7571 F.removeFnAttr(FnAttrName);
7572 F.addFnAttr(FnAttrName);
7578 Triple T(M.getTargetTriple());
7579 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7582 uint64_t BTEValue = 0;
7583 uint64_t BPPLRValue = 0;
7584 uint64_t GCSValue = 0;
7585 uint64_t SRAValue = 0;
7586 uint64_t SRAALLValue = 0;
7587 uint64_t SRABKeyValue = 0;
7589 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7593 if (
Op->getNumOperands() != 3)
7602 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7603 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7604 : IDStr ==
"guarded-control-stack" ? &GCSValue
7605 : IDStr ==
"sign-return-address" ? &SRAValue
7606 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7607 : IDStr ==
"sign-return-address-with-bkey"
7613 *ValPtr = CI->getZExtValue();
7619 bool BTE = BTEValue == 1;
7620 bool BPPLR = BPPLRValue == 1;
7621 bool GCS = GCSValue == 1;
7622 bool SRA = SRAValue == 1;
7625 if (SRA && SRAALLValue == 1)
7626 SignTypeValue =
"all";
7629 if (SRA && SRABKeyValue == 1)
7630 SignKeyValue =
"b_key";
7632 for (
Function &
F : M.getFunctionList()) {
7633 if (
F.isDeclaration())
7640 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7641 A.isValid() &&
"none" ==
A.getValueAsString()) {
7642 F.removeFnAttr(
"sign-return-address");
7643 F.removeFnAttr(
"sign-return-address-key");
7659 if (SRAALLValue == 1)
7661 if (SRABKeyValue == 1)
7688 if (
T->getNumOperands() < 1)
7693 if (S->getString().starts_with(
"llvm.vectorizer."))
7699 StringRef OldPrefix =
"llvm.vectorizer.";
7702 if (OldTag ==
"llvm.vectorizer.unroll")
7714 if (
T->getNumOperands() < 1)
7726 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7739 Ops.reserve(
T->getNumOperands());
7740 Ops.push_back(NewTag);
7741 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7742 Ops.push_back(
T->getOperand(
I));
7759 if (
T->isDistinct()) {
7760 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7772 Ops.reserve(
T->getNumOperands());
7783 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7784 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7785 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7788 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7790 auto I =
DL.find(
"-n64-");
7792 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7797 std::string Res =
DL.str();
7800 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7801 Res.append(Res.empty() ?
"G1" :
"-G1");
7809 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7810 Res.append(
"-ni:7:8:9");
7812 if (
DL.ends_with(
"ni:7"))
7814 if (
DL.ends_with(
"ni:7:8"))
7819 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7820 Res.append(
"-p7:160:256:256:32");
7821 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7822 Res.append(
"-p8:128:128:128:48");
7823 constexpr StringRef OldP8(
"-p8:128:128-");
7824 if (
DL.contains(OldP8))
7825 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7826 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7827 Res.append(
"-p9:192:256:256:32");
7832 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7833 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7834 Res.append((
"-" + AS +
":32:32").str());
7839 if (!
DL.contains(
"m:e"))
7840 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7845 if (
T.isSystemZ() && !
DL.empty()) {
7847 if (!
DL.contains(
"-S64"))
7848 return "E-S64" +
DL.drop_front(1).str();
7852 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7855 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7856 if (!
DL.contains(AddrSpaces)) {
7858 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7859 if (R.match(Res, &
Groups))
7865 if (
T.isAArch64()) {
7867 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7868 Res.append(
"-Fn32");
7869 AddPtr32Ptr64AddrSpaces();
7873 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7877 std::string I64 =
"-i64:64";
7878 std::string I128 =
"-i128:128";
7880 size_t Pos = Res.find(I64);
7881 if (Pos !=
size_t(-1))
7882 Res.insert(Pos + I64.size(), I128);
7886 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7887 size_t Pos = Res.find(
"-S128");
7890 Res.insert(Pos,
"-f64:32:64");
7896 AddPtr32Ptr64AddrSpaces();
7904 if (!
T.isOSIAMCU()) {
7905 std::string I128 =
"-i128:128";
7908 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7909 if (R.match(Res, &
Groups))
7917 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7919 auto I =
Ref.find(
"-f80:32-");
7921 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7929 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7932 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7933 B.removeAttribute(
"no-frame-pointer-elim");
7935 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7937 if (FramePointer !=
"all")
7938 FramePointer =
"non-leaf";
7939 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7941 if (!FramePointer.
empty())
7942 B.addAttribute(
"frame-pointer", FramePointer);
7944 A =
B.getAttribute(
"null-pointer-is-valid");
7947 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7948 B.removeAttribute(
"null-pointer-is-valid");
7949 if (NullPointerIsValid)
7950 B.addAttribute(Attribute::NullPointerIsValid);
7953 A =
B.getAttribute(
"uniform-work-group-size");
7957 bool IsTrue = Val ==
"true";
7958 B.removeAttribute(
"uniform-work-group-size");
7960 B.addAttribute(
"uniform-work-group-size");
7971 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)
static unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID, SmallVectorImpl< Type * > &OverloadTys)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFAddUpgrade(StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static cl::opt< bool > DisableAutoUpgradeDebugInfo("disable-auto-upgrade-debug-info", cl::desc("Disable autoupgrade of debug info"))
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F, StringRef Name)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_CASE(ID_SUFFIX, NAME)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static FCmpInst::Predicate getVPFPPredicateFromMD(const Value *Op)
static GCRegistry::Add< ShadowStackGC > C("shadow-stack", "Very portable GC for uncooperative code generators")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
This file contains constants used for implementing Dwarf debug support.
Module.h This file contains the declarations for the Module class.
const AbstractManglingParser< Derived, Alloc >::OperatorInfo AbstractManglingParser< Derived, Alloc >::Ops[]
static bool isZero(Value *V, const DataLayout &DL, DominatorTree *DT, AssumptionCache *AC)
NVPTX address space definition.
This file contains the definitions of the enumerations and flags associated with NVVM Intrinsics,...
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
static SymbolRef::Type getType(const Symbol *Sym)
LocallyHashedType DenseMapInfo< LocallyHashedType >::Empty
static const X86InstrFMA3Group Groups[]
Class for arbitrary precision integers.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
Class to represent array types.
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Type * getElementType() const
an instruction that atomically reads a memory location, combines it with another value,...
void setVolatile(bool V)
Specify whether this is a volatile RMW or not.
BinOp
This enumeration lists the possible modifications atomicrmw can make.
@ USubCond
Subtract only if no unsigned overflow.
@ Min
*p = old <signed v ? old : v
@ USubSat
*p = usub.sat(old, v) usub.sat matches the behavior of llvm.usub.sat.
@ UIncWrap
Increment one up to a maximum value.
@ Max
*p = old >signed v ? old : v
@ FMin
*p = minnum(old, v) minnum matches the behavior of llvm.minnum.
@ FMax
*p = maxnum(old, v) maxnum matches the behavior of llvm.maxnum.
@ UDecWrap
Decrement one until a minimum value or zero.
bool isFloatingPointOperation() const
This class stores enough information to efficiently remove some attributes from an existing AttrBuild...
AttributeMask & addAttribute(Attribute::AttrKind Val)
Add an attribute to the mask.
Functions, function parameters, and return types can have attributes to indicate how they should be t...
static LLVM_ABI Attribute getWithStackAlignment(LLVMContext &Context, Align Alignment)
static LLVM_ABI Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
void setCallingConv(CallingConv::ID CC)
LLVM_ABI void getOperandBundlesAsDefs(SmallVectorImpl< OperandBundleDef > &Defs) const
Return the list of operand bundles attached to this instruction as a vector of OperandBundleDefs.
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
CallingConv::ID getCallingConv() const
Value * getCalledOperand() const
void setAttributes(AttributeList A)
Set the attributes for this call.
Value * getArgOperand(unsigned i) const
FunctionType * getFunctionType() const
LLVM_ABI Intrinsic::ID getIntrinsicID() const
Returns the intrinsic ID of the intrinsic called or Intrinsic::not_intrinsic if the called function i...
iterator_range< User::op_iterator > args()
Iteration adapter for range-for loops.
void setCalledOperand(Value *V)
unsigned arg_size() const
AttributeList getAttributes() const
Return the attributes for this call.
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
This class represents a function call, abstracting a target machine's calling convention.
void setTailCallKind(TailCallKind TCK)
static LLVM_ABI CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", InsertPosition InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
static LLVM_ABI bool castIsValid(Instruction::CastOps op, Type *SrcTy, Type *DstTy)
This method can be used to determine if a cast from SrcTy to DstTy using Opcode op is valid or not.
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
unsigned getNumParams() const
Return the number of fixed parameters this function type requires.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
StringSwitch & StartsWith(StringLiteral S, T Value)
StringSwitch & Cases(std::initializer_list< StringLiteral > CaseStrings, T Value)
Class to represent struct types.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
unsigned getNumElements() const
Random access to the elements.
Type * getElementType(unsigned N) const
The TimeTraceScope is a helper class to call the begin and end functions of the time trace profiler.
Triple - Helper class for working with autoconf configuration names.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)
bool isVectorTy() const
True if this is an instance of VectorType.
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
LLVM_ABI unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)
LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isPtrOrPtrVectorTy() const
Return true if this is a pointer type or a vector of pointer types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isFPOrFPVectorTy() const
Return true if this is a FP type or a vector of FP.
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
static LLVM_ABI Type * getBFloatTy(LLVMContext &C)
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
A Use represents the edge between a Value definition and its users.
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
LLVM_ABI void print(raw_ostream &O, bool IsForDebug=false) const
Implement operator<< on Value.
LLVM_ABI void setName(const Twine &Name)
Change the name of the value.
LLVM_ABI void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
LLVMContext & getContext() const
All values hold a context through their type.
iterator_range< user_iterator > users()
LLVM_ABI const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
LLVM_ABI void takeName(Value *V)
Transfer the name from V to this value.
Base class of all SIMD vector types.
static VectorType * getInteger(VectorType *VTy)
This static method gets a VectorType with the same number of elements as the input type,...
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
constexpr ScalarTy getFixedValue() const
const ParentTy * getParent() const
self_iterator getIterator()
A raw_ostream that writes to an SmallVector or SmallString.
StringRef str() const
Return a StringRef for the vector contents.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ LOCAL_ADDRESS
Address space for local memory.
@ FLAT_ADDRESS
Address space for flat memory.
@ PRIVATE_ADDRESS
Address space for private memory.
@ PTX_Kernel
Call to a PTX kernel. Passes all arguments in parameter space.
std::optional< ABIType > parseABIType(StringRef S)
Parse the string spelling used by the "float-abi" IR module flag into an ABIType.
LLVM_ABI std::optional< Function * > remangleIntrinsicFunction(Function *F)
LLVM_ABI Function * getOrInsertDeclaration(Module *M, ID id, ArrayRef< Type * > OverloadTys={})
Look up the Function declaration of the intrinsic id in the Module M.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
bool is_contained(R &&Range, const E &Element)
Returns true if Element is found in Range.
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
@ Default
The result value is uniform if and only if all operands are uniform.
LLVM_ABI bool verifyModule(const Module &M, raw_ostream *OS=nullptr, bool *BrokenDebugInfo=nullptr)
Check a module for errors.
LLVM_ABI void reportFatalUsageError(Error Err)
Report a fatal error that does not indicate a bug in LLVM.
void swap(llvm::BitVector &LHS, llvm::BitVector &RHS)
Implement std::swap in terms of BitVector swap.
This struct is a compact representation of a valid (non-zero power of two) alignment.
Represents the full denormal controls for a function, including the default mode and the f32 specific...
Represent subnormal handling kind for floating point instruction inputs and outputs.
static constexpr DenormalMode getInvalid()
constexpr bool isValid() const
static constexpr DenormalMode getIEEE()
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.