35#include "llvm/IR/IntrinsicsAArch64.h"
36#include "llvm/IR/IntrinsicsAMDGPU.h"
37#include "llvm/IR/IntrinsicsARM.h"
38#include "llvm/IR/IntrinsicsNVPTX.h"
39#include "llvm/IR/IntrinsicsRISCV.h"
40#include "llvm/IR/IntrinsicsWebAssembly.h"
41#include "llvm/IR/IntrinsicsX86.h"
66 cl::desc(
"Disable autoupgrade of debug info"));
85 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
100 Type *LastArgType =
F->getFunctionType()->getParamType(
101 F->getFunctionType()->getNumParams() - 1);
116 if (
F->getReturnType()->isVectorTy())
129 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
130 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
147 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
148 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
162 if (
F->getReturnType()->getScalarType()->isBFloatTy())
172 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
186 if (Name.consume_front(
"avx."))
187 return (Name.starts_with(
"blend.p") ||
188 Name ==
"cvt.ps2.pd.256" ||
189 Name ==
"cvtdq2.pd.256" ||
190 Name ==
"cvtdq2.ps.256" ||
191 Name.starts_with(
"movnt.") ||
192 Name.starts_with(
"sqrt.p") ||
193 Name.starts_with(
"storeu.") ||
194 Name.starts_with(
"vbroadcast.s") ||
195 Name.starts_with(
"vbroadcastf128") ||
196 Name.starts_with(
"vextractf128.") ||
197 Name.starts_with(
"vinsertf128.") ||
198 Name.starts_with(
"vperm2f128.") ||
199 Name.starts_with(
"vpermil."));
201 if (Name.consume_front(
"avx2."))
202 return (Name ==
"movntdqa" ||
203 Name.starts_with(
"pabs.") ||
204 Name.starts_with(
"padds.") ||
205 Name.starts_with(
"paddus.") ||
206 Name.starts_with(
"pblendd.") ||
208 Name.starts_with(
"pbroadcast") ||
209 Name.starts_with(
"pcmpeq.") ||
210 Name.starts_with(
"pcmpgt.") ||
211 Name.starts_with(
"pmax") ||
212 Name.starts_with(
"pmin") ||
213 Name.starts_with(
"pmovsx") ||
214 Name.starts_with(
"pmovzx") ||
216 Name ==
"pmulu.dq" ||
217 Name.starts_with(
"psll.dq") ||
218 Name.starts_with(
"psrl.dq") ||
219 Name.starts_with(
"psubs.") ||
220 Name.starts_with(
"psubus.") ||
221 Name.starts_with(
"vbroadcast") ||
222 Name ==
"vbroadcasti128" ||
223 Name ==
"vextracti128" ||
224 Name ==
"vinserti128" ||
225 Name ==
"vperm2i128");
227 if (Name.consume_front(
"avx512.")) {
228 if (Name.consume_front(
"mask."))
230 return (Name.starts_with(
"add.p") ||
231 Name.starts_with(
"and.") ||
232 Name.starts_with(
"andn.") ||
233 Name.starts_with(
"broadcast.s") ||
234 Name.starts_with(
"broadcastf32x4.") ||
235 Name.starts_with(
"broadcastf32x8.") ||
236 Name.starts_with(
"broadcastf64x2.") ||
237 Name.starts_with(
"broadcastf64x4.") ||
238 Name.starts_with(
"broadcasti32x4.") ||
239 Name.starts_with(
"broadcasti32x8.") ||
240 Name.starts_with(
"broadcasti64x2.") ||
241 Name.starts_with(
"broadcasti64x4.") ||
242 Name.starts_with(
"cmp.b") ||
243 Name.starts_with(
"cmp.d") ||
244 Name.starts_with(
"cmp.q") ||
245 Name.starts_with(
"cmp.w") ||
246 Name.starts_with(
"compress.b") ||
247 Name.starts_with(
"compress.d") ||
248 Name.starts_with(
"compress.p") ||
249 Name.starts_with(
"compress.q") ||
250 Name.starts_with(
"compress.store.") ||
251 Name.starts_with(
"compress.w") ||
252 Name.starts_with(
"conflict.") ||
253 Name.starts_with(
"cvtdq2pd.") ||
254 Name.starts_with(
"cvtdq2ps.") ||
255 Name ==
"cvtpd2dq.256" ||
256 Name ==
"cvtpd2ps.256" ||
257 Name ==
"cvtps2pd.128" ||
258 Name ==
"cvtps2pd.256" ||
259 Name.starts_with(
"cvtqq2pd.") ||
260 Name ==
"cvtqq2ps.256" ||
261 Name ==
"cvtqq2ps.512" ||
262 Name ==
"cvttpd2dq.256" ||
263 Name ==
"cvttps2dq.128" ||
264 Name ==
"cvttps2dq.256" ||
265 Name.starts_with(
"cvtudq2pd.") ||
266 Name.starts_with(
"cvtudq2ps.") ||
267 Name.starts_with(
"cvtuqq2pd.") ||
268 Name ==
"cvtuqq2ps.256" ||
269 Name ==
"cvtuqq2ps.512" ||
270 Name.starts_with(
"dbpsadbw.") ||
271 Name.starts_with(
"div.p") ||
272 Name.starts_with(
"expand.b") ||
273 Name.starts_with(
"expand.d") ||
274 Name.starts_with(
"expand.load.") ||
275 Name.starts_with(
"expand.p") ||
276 Name.starts_with(
"expand.q") ||
277 Name.starts_with(
"expand.w") ||
278 Name.starts_with(
"fpclass.p") ||
279 Name.starts_with(
"insert") ||
280 Name.starts_with(
"load.") ||
281 Name.starts_with(
"loadu.") ||
282 Name.starts_with(
"lzcnt.") ||
283 Name.starts_with(
"max.p") ||
284 Name.starts_with(
"min.p") ||
285 Name.starts_with(
"movddup") ||
286 Name.starts_with(
"move.s") ||
287 Name.starts_with(
"movshdup") ||
288 Name.starts_with(
"movsldup") ||
289 Name.starts_with(
"mul.p") ||
290 Name.starts_with(
"or.") ||
291 Name.starts_with(
"pabs.") ||
292 Name.starts_with(
"packssdw.") ||
293 Name.starts_with(
"packsswb.") ||
294 Name.starts_with(
"packusdw.") ||
295 Name.starts_with(
"packuswb.") ||
296 Name.starts_with(
"padd.") ||
297 Name.starts_with(
"padds.") ||
298 Name.starts_with(
"paddus.") ||
299 Name.starts_with(
"palignr.") ||
300 Name.starts_with(
"pand.") ||
301 Name.starts_with(
"pandn.") ||
302 Name.starts_with(
"pavg") ||
303 Name.starts_with(
"pbroadcast") ||
304 Name.starts_with(
"pcmpeq.") ||
305 Name.starts_with(
"pcmpgt.") ||
306 Name.starts_with(
"perm.df.") ||
307 Name.starts_with(
"perm.di.") ||
308 Name.starts_with(
"permvar.") ||
309 Name.starts_with(
"pmaddubs.w.") ||
310 Name.starts_with(
"pmaddw.d.") ||
311 Name.starts_with(
"pmax") ||
312 Name.starts_with(
"pmin") ||
313 Name ==
"pmov.qd.256" ||
314 Name ==
"pmov.qd.512" ||
315 Name ==
"pmov.wb.256" ||
316 Name ==
"pmov.wb.512" ||
317 Name.starts_with(
"pmovsx") ||
318 Name.starts_with(
"pmovzx") ||
319 Name.starts_with(
"pmul.dq.") ||
320 Name.starts_with(
"pmul.hr.sw.") ||
321 Name.starts_with(
"pmulh.w.") ||
322 Name.starts_with(
"pmulhu.w.") ||
323 Name.starts_with(
"pmull.") ||
324 Name.starts_with(
"pmultishift.qb.") ||
325 Name.starts_with(
"pmulu.dq.") ||
326 Name.starts_with(
"por.") ||
327 Name.starts_with(
"prol.") ||
328 Name.starts_with(
"prolv.") ||
329 Name.starts_with(
"pror.") ||
330 Name.starts_with(
"prorv.") ||
331 Name.starts_with(
"pshuf.b.") ||
332 Name.starts_with(
"pshuf.d.") ||
333 Name.starts_with(
"pshufh.w.") ||
334 Name.starts_with(
"pshufl.w.") ||
335 Name.starts_with(
"psll.d") ||
336 Name.starts_with(
"psll.q") ||
337 Name.starts_with(
"psll.w") ||
338 Name.starts_with(
"pslli") ||
339 Name.starts_with(
"psllv") ||
340 Name.starts_with(
"psra.d") ||
341 Name.starts_with(
"psra.q") ||
342 Name.starts_with(
"psra.w") ||
343 Name.starts_with(
"psrai") ||
344 Name.starts_with(
"psrav") ||
345 Name.starts_with(
"psrl.d") ||
346 Name.starts_with(
"psrl.q") ||
347 Name.starts_with(
"psrl.w") ||
348 Name.starts_with(
"psrli") ||
349 Name.starts_with(
"psrlv") ||
350 Name.starts_with(
"psub.") ||
351 Name.starts_with(
"psubs.") ||
352 Name.starts_with(
"psubus.") ||
353 Name.starts_with(
"pternlog.") ||
354 Name.starts_with(
"punpckh") ||
355 Name.starts_with(
"punpckl") ||
356 Name.starts_with(
"pxor.") ||
357 Name.starts_with(
"shuf.f") ||
358 Name.starts_with(
"shuf.i") ||
359 Name.starts_with(
"shuf.p") ||
360 Name.starts_with(
"sqrt.p") ||
361 Name.starts_with(
"store.b.") ||
362 Name.starts_with(
"store.d.") ||
363 Name.starts_with(
"store.p") ||
364 Name.starts_with(
"store.q.") ||
365 Name.starts_with(
"store.w.") ||
366 Name ==
"store.ss" ||
367 Name.starts_with(
"storeu.") ||
368 Name.starts_with(
"sub.p") ||
369 Name.starts_with(
"ucmp.") ||
370 Name.starts_with(
"unpckh.") ||
371 Name.starts_with(
"unpckl.") ||
372 Name.starts_with(
"valign.") ||
373 Name ==
"vcvtph2ps.128" ||
374 Name ==
"vcvtph2ps.256" ||
375 Name.starts_with(
"vextract") ||
376 Name.starts_with(
"vfmadd.") ||
377 Name.starts_with(
"vfmaddsub.") ||
378 Name.starts_with(
"vfnmadd.") ||
379 Name.starts_with(
"vfnmsub.") ||
380 Name.starts_with(
"vpdpbusd.") ||
381 Name.starts_with(
"vpdpbusds.") ||
382 Name.starts_with(
"vpdpwssd.") ||
383 Name.starts_with(
"vpdpwssds.") ||
384 Name.starts_with(
"vpermi2var.") ||
385 Name.starts_with(
"vpermil.p") ||
386 Name.starts_with(
"vpermilvar.") ||
387 Name.starts_with(
"vpermt2var.") ||
388 Name.starts_with(
"vpmadd52") ||
389 Name.starts_with(
"vpshld.") ||
390 Name.starts_with(
"vpshldv.") ||
391 Name.starts_with(
"vpshrd.") ||
392 Name.starts_with(
"vpshrdv.") ||
393 Name.starts_with(
"vpshufbitqmb.") ||
394 Name.starts_with(
"xor."));
396 if (Name.consume_front(
"mask3."))
398 return (Name.starts_with(
"vfmadd.") ||
399 Name.starts_with(
"vfmaddsub.") ||
400 Name.starts_with(
"vfmsub.") ||
401 Name.starts_with(
"vfmsubadd.") ||
402 Name.starts_with(
"vfnmsub."));
404 if (Name.consume_front(
"maskz."))
406 return (Name.starts_with(
"pternlog.") ||
407 Name.starts_with(
"vfmadd.") ||
408 Name.starts_with(
"vfmaddsub.") ||
409 Name.starts_with(
"vpdpbusd.") ||
410 Name.starts_with(
"vpdpbusds.") ||
411 Name.starts_with(
"vpdpwssd.") ||
412 Name.starts_with(
"vpdpwssds.") ||
413 Name.starts_with(
"vpermt2var.") ||
414 Name.starts_with(
"vpmadd52") ||
415 Name.starts_with(
"vpshldv.") ||
416 Name.starts_with(
"vpshrdv."));
419 return (Name ==
"movntdqa" ||
420 Name ==
"pmul.dq.512" ||
421 Name ==
"pmulu.dq.512" ||
422 Name.starts_with(
"broadcastm") ||
423 Name.starts_with(
"cmp.p") ||
424 Name.starts_with(
"cvtb2mask.") ||
425 Name.starts_with(
"cvtd2mask.") ||
426 Name.starts_with(
"cvtmask2") ||
427 Name.starts_with(
"cvtq2mask.") ||
428 Name ==
"cvtusi2sd" ||
429 Name.starts_with(
"cvtw2mask.") ||
434 Name ==
"kortestc.w" ||
435 Name ==
"kortestz.w" ||
436 Name.starts_with(
"kunpck") ||
439 Name.starts_with(
"padds.") ||
440 Name.starts_with(
"pbroadcast") ||
441 Name.starts_with(
"prol") ||
442 Name.starts_with(
"pror") ||
443 Name.starts_with(
"psll.dq") ||
444 Name.starts_with(
"psrl.dq") ||
445 Name.starts_with(
"psubs.") ||
446 Name.starts_with(
"ptestm") ||
447 Name.starts_with(
"ptestnm") ||
448 Name.starts_with(
"storent.") ||
449 Name.starts_with(
"vbroadcast.s") ||
450 Name.starts_with(
"vpshld.") ||
451 Name.starts_with(
"vpshrd."));
454 if (Name.consume_front(
"fma."))
455 return (Name.starts_with(
"vfmadd.") ||
456 Name.starts_with(
"vfmsub.") ||
457 Name.starts_with(
"vfmsubadd.") ||
458 Name.starts_with(
"vfnmadd.") ||
459 Name.starts_with(
"vfnmsub."));
461 if (Name.consume_front(
"fma4."))
462 return Name.starts_with(
"vfmadd.s");
464 if (Name.consume_front(
"sse."))
465 return (Name ==
"add.ss" ||
466 Name ==
"cvtsi2ss" ||
467 Name ==
"cvtsi642ss" ||
470 Name.starts_with(
"sqrt.p") ||
472 Name.starts_with(
"storeu.") ||
475 if (Name.consume_front(
"sse2."))
476 return (Name ==
"add.sd" ||
477 Name ==
"cvtdq2pd" ||
478 Name ==
"cvtdq2ps" ||
479 Name ==
"cvtps2pd" ||
480 Name ==
"cvtsi2sd" ||
481 Name ==
"cvtsi642sd" ||
482 Name ==
"cvtss2sd" ||
485 Name.starts_with(
"padds.") ||
486 Name.starts_with(
"paddus.") ||
487 Name.starts_with(
"pcmpeq.") ||
488 Name.starts_with(
"pcmpgt.") ||
493 Name ==
"pmulu.dq" ||
494 Name.starts_with(
"pshuf") ||
495 Name.starts_with(
"psll.dq") ||
496 Name.starts_with(
"psrl.dq") ||
497 Name.starts_with(
"psubs.") ||
498 Name.starts_with(
"psubus.") ||
499 Name.starts_with(
"sqrt.p") ||
501 Name ==
"storel.dq" ||
502 Name.starts_with(
"storeu.") ||
505 if (Name.consume_front(
"sse41."))
506 return (Name.starts_with(
"blendp") ||
507 Name ==
"movntdqa" ||
517 Name.starts_with(
"pmovsx") ||
518 Name.starts_with(
"pmovzx") ||
521 if (Name.consume_front(
"sse42."))
522 return Name ==
"crc32.64.8";
524 if (Name.consume_front(
"sse4a."))
525 return Name.starts_with(
"movnt.");
527 if (Name.consume_front(
"ssse3."))
528 return (Name ==
"pabs.b.128" ||
529 Name ==
"pabs.d.128" ||
530 Name ==
"pabs.w.128");
532 if (Name.consume_front(
"xop."))
533 return (Name ==
"vpcmov" ||
534 Name ==
"vpcmov.256" ||
535 Name.starts_with(
"vpcom") ||
536 Name.starts_with(
"vprot"));
538 if (Name.consume_front(
"bmi."))
539 return (Name.starts_with(
"pdep.") ||
540 Name.starts_with(
"pext."));
542 return (Name ==
"addcarry.u32" ||
543 Name ==
"addcarry.u64" ||
544 Name ==
"addcarryx.u32" ||
545 Name ==
"addcarryx.u64" ||
546 Name ==
"subborrow.u32" ||
547 Name ==
"subborrow.u64" ||
548 Name.starts_with(
"vcvtph2ps."));
554 if (!Name.consume_front(
"x86."))
562 if (Name ==
"rdtscp") {
564 if (
F->getFunctionType()->getNumParams() == 0)
569 Intrinsic::x86_rdtscp);
576 if (Name.consume_front(
"sse41.ptest")) {
578 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
579 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
580 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
593 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
594 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
595 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
596 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
597 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
598 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
603 if (Name.consume_front(
"avx512.")) {
604 if (Name.consume_front(
"mask.cmp.")) {
607 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
608 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
609 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
610 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
611 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
612 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
616 }
else if (Name.starts_with(
"vpdpbusd.") ||
617 Name.starts_with(
"vpdpbusds.")) {
620 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
621 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
622 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
623 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
624 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
625 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
629 }
else if (Name.starts_with(
"vpdpwssd.") ||
630 Name.starts_with(
"vpdpwssds.")) {
633 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
634 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
635 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
636 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
637 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
638 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
646 if (Name.consume_front(
"avx2.")) {
647 if (Name.consume_front(
"vpdpb")) {
650 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
651 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
652 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
653 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
654 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
655 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
656 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
657 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
658 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
659 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
660 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
661 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
665 }
else if (Name.consume_front(
"vpdpw")) {
668 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
669 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
670 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
671 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
672 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
673 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
674 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
675 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
676 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
677 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
678 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
679 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
687 if (Name.consume_front(
"avx10.")) {
688 if (Name.consume_front(
"vpdpb")) {
691 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
692 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
693 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
694 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
695 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
696 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
700 }
else if (Name.consume_front(
"vpdpw")) {
702 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
703 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
704 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
705 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
706 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
707 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
715 if (Name.consume_front(
"avx512bf16.")) {
718 .
Case(
"cvtne2ps2bf16.128",
719 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
720 .
Case(
"cvtne2ps2bf16.256",
721 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
722 .
Case(
"cvtne2ps2bf16.512",
723 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
724 .
Case(
"mask.cvtneps2bf16.128",
725 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
726 .
Case(
"cvtneps2bf16.256",
727 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
728 .
Case(
"cvtneps2bf16.512",
729 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
736 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
737 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
738 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
745 if (Name.consume_front(
"xop.")) {
747 if (Name.starts_with(
"vpermil2")) {
750 auto Idx =
F->getFunctionType()->getParamType(2);
751 if (Idx->isFPOrFPVectorTy()) {
752 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
753 unsigned EltSize = Idx->getScalarSizeInBits();
754 if (EltSize == 64 && IdxSize == 128)
755 ID = Intrinsic::x86_xop_vpermil2pd;
756 else if (EltSize == 32 && IdxSize == 128)
757 ID = Intrinsic::x86_xop_vpermil2ps;
758 else if (EltSize == 64 && IdxSize == 256)
759 ID = Intrinsic::x86_xop_vpermil2pd_256;
761 ID = Intrinsic::x86_xop_vpermil2ps_256;
763 }
else if (
F->arg_size() == 2)
766 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
767 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
778 if (Name ==
"seh.recoverfp") {
780 Intrinsic::eh_recoverfp);
792 if (Name.starts_with(
"rbit")) {
795 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
799 if (Name ==
"thread.pointer") {
802 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
806 bool Neon = Name.consume_front(
"neon.");
811 if (Name.consume_front(
"bfdot.")) {
815 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
820 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
821 assert((OperandWidth == 64 || OperandWidth == 128) &&
822 "Unexpected operand width");
824 std::array<Type *, 2> Tys{
835 if (Name.consume_front(
"bfm")) {
837 if (Name.consume_back(
".v4f32.v16i8")) {
883 F->arg_begin()->getType());
887 if (Name.consume_front(
"vst")) {
889 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
893 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
894 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
897 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
898 Intrinsic::arm_neon_vst4lane};
900 auto fArgs =
F->getFunctionType()->params();
901 Type *Tys[] = {fArgs[0], fArgs[1]};
904 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
907 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
916 if (Name.consume_front(
"mve.")) {
918 if (Name ==
"vctp64") {
928 if (Name.starts_with(
"vrintn.v")) {
930 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
935 if (Name.consume_back(
".v4i1")) {
937 if (Name.consume_back(
".predicated.v2i64.v4i32"))
939 return Name ==
"mull.int" || Name ==
"vqdmull";
941 if (Name.consume_back(
".v2i64")) {
943 bool IsGather = Name.consume_front(
"vldr.gather.");
944 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
945 if (Name.consume_front(
"base.")) {
947 Name.consume_front(
"wb.");
950 return Name ==
"predicated.v2i64";
953 if (Name.consume_front(
"offset.predicated."))
954 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
955 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
968 if (Name.consume_front(
"cde.vcx")) {
970 if (Name.consume_back(
".predicated.v2i64.v4i1"))
972 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
973 Name ==
"3q" || Name ==
"3qa";
987 F->arg_begin()->getType());
991 if (Name.starts_with(
"addp")) {
993 if (
F->arg_size() != 2)
996 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
998 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1004 if (Name.starts_with(
"bfcvt")) {
1010 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1017 if (Name.consume_front(
"sve.")) {
1019 if (Name.consume_front(
"bf")) {
1020 if (Name ==
"mmla") {
1021 Type *Tys[] = {
F->getReturnType(),
1022 std::next(
F->arg_begin())->getType()};
1024 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1027 if (Name.consume_back(
".lane")) {
1031 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1032 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1033 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1045 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1050 if (Name.consume_front(
"addqv")) {
1052 if (!
F->getReturnType()->isFPOrFPVectorTy())
1055 auto Args =
F->getFunctionType()->params();
1056 Type *Tys[] = {
F->getReturnType(), Args[1]};
1058 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1062 if (Name.consume_front(
"ld")) {
1064 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1065 if (LdRegex.
match(Name)) {
1071 "Expected 2 arguments for ld* intrinsic.");
1072 Type *PtrTy =
F->getArg(1)->getType();
1075 Intrinsic::aarch64_sve_ld2_sret,
1076 Intrinsic::aarch64_sve_ld3_sret,
1077 Intrinsic::aarch64_sve_ld4_sret,
1080 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1086 if (Name.consume_front(
"tuple.")) {
1088 if (Name.starts_with(
"get")) {
1090 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1092 F->getParent(), Intrinsic::vector_extract, Tys);
1096 if (Name.starts_with(
"set")) {
1098 auto Args =
F->getFunctionType()->params();
1099 Type *Tys[] = {Args[0], Args[2], Args[1]};
1101 F->getParent(), Intrinsic::vector_insert, Tys);
1105 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1106 if (CreateTupleRegex.
match(Name)) {
1108 auto Args =
F->getFunctionType()->params();
1109 Type *Tys[] = {
F->getReturnType(), Args[1]};
1111 F->getParent(), Intrinsic::vector_insert, Tys);
1117 if (Name.starts_with(
"rev.nxv")) {
1120 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1126 if (Name.consume_front(
"sme.")) {
1128 if (Name.consume_front(
"ftmopa.")) {
1133 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1134 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1151 if (Name.consume_front(
"cp.async.bulk.tensor.g2s.")) {
1155 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d)
1157 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d)
1159 Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d)
1160 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d)
1161 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d)
1162 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d)
1163 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d)
1164 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d)
1173 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1187 size_t FlagStartIndex =
F->getFunctionType()->getNumParams() - 3;
1188 Type *ArgType =
F->getFunctionType()->getParamType(FlagStartIndex);
1213 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1216 auto [RedOpName, ShapeName] = Name.split(
'.');
1221 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1222 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1223 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1224 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1225 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1226 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1227 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1228 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1234 if (Name.consume_front(
"mapa.shared.cluster"))
1235 if (
F->getReturnType()->getPointerAddressSpace() ==
1237 return Intrinsic::nvvm_mapa_shared_cluster;
1239 if (Name.consume_front(
"cp.async.bulk.")) {
1242 .
Case(
"global.to.shared.cluster",
1243 Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster)
1244 .
Case(
"shared.cta.to.cluster",
1245 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1249 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1259 if (!Name.consume_front(
"tcgen05.commit."))
1262 if (Name.consume_front(
"shared."))
1264 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1265 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1268 if (Name.consume_front(
"mc.shared.")) {
1270 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1274 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1275 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1283 if (Name.consume_front(
"fma.rn."))
1285 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1286 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1287 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1288 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1291 if (Name.consume_front(
"fmax."))
1293 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1294 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1295 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1296 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1297 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1298 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1299 .
Case(
"ftz.nan.xorsign.abs.bf16",
1300 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1301 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1302 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1303 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1304 .
Case(
"ftz.xorsign.abs.bf16x2",
1305 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1306 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1307 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1308 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1309 .
Case(
"nan.xorsign.abs.bf16x2",
1310 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1311 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1312 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1315 if (Name.consume_front(
"fmin."))
1317 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1318 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1319 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1320 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1321 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1322 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1323 .
Case(
"ftz.nan.xorsign.abs.bf16",
1324 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1325 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1326 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1327 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1328 .
Case(
"ftz.xorsign.abs.bf16x2",
1329 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1330 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1331 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1332 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1333 .
Case(
"nan.xorsign.abs.bf16x2",
1334 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1335 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1336 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1339 if (Name.consume_front(
"neg."))
1341 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1342 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1350 if (!Name.consume_front(
"tcgen05.mma."))
1354 if (Name.starts_with(
"ws"))
1357 return F->getIntrinsicID();
1361 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1362 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1363 Name.consume_front(
"param");
1369 if (Name.starts_with(
"to.fp16")) {
1373 FuncTy->getReturnType());
1376 if (Name.starts_with(
"from.fp16")) {
1380 FuncTy->getReturnType());
1392 if (Defaults.empty())
1404 if (
F->arg_size() >= FullDecl->
arg_size())
1409 if (
F->arg_size() < FirstDefault)
1417 bool CanUpgradeDebugIntrinsicsToRecords) {
1418 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1423 if (!Name.consume_front(
"llvm.") || Name.empty())
1429 bool IsArm = Name.consume_front(
"arm.");
1430 if (IsArm || Name.consume_front(
"aarch64.")) {
1436 if (Name.consume_front(
"amdgcn.")) {
1437 if (Name ==
"alignbit") {
1440 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1444 if (Name.consume_front(
"atomic.")) {
1445 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1446 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1455 switch (
F->getIntrinsicID()) {
1459 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1460 if (
F->arg_size() == 7) {
1465 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1466 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1467 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1468 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1469 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1470 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1471 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1472 if (
F->arg_size() == 8) {
1479 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1480 Name.consume_front(
"flat.atomic.")) {
1481 if (Name.starts_with(
"fadd") ||
1483 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1484 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1492 if (Name.starts_with(
"ldexp.")) {
1495 F->getParent(), Intrinsic::ldexp,
1496 {F->getReturnType(), F->getArg(1)->getType()});
1505 if (
F->arg_size() == 1) {
1506 if (Name.consume_front(
"convert.")) {
1520 F->arg_begin()->getType());
1526 if (Name ==
"coro.end" &&
1527 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1528 CoroEndID = Intrinsic::coro_end;
1529 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1530 CoroEndID = Intrinsic::coro_end_async;
1541 if (Name.consume_front(
"dbg.")) {
1543 if (CanUpgradeDebugIntrinsicsToRecords) {
1544 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1545 Name ==
"declare" || Name ==
"label") {
1554 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1557 Intrinsic::dbg_value);
1564 if (Name.consume_front(
"experimental.vector.")) {
1570 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1571 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1572 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1573 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1574 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1576 Intrinsic::vector_partial_reduce_add)
1579 const auto *FT =
F->getFunctionType();
1581 if (ID == Intrinsic::vector_extract ||
1582 ID == Intrinsic::vector_interleave2)
1585 if (ID != Intrinsic::vector_interleave2)
1587 if (ID == Intrinsic::vector_insert ||
1588 ID == Intrinsic::vector_partial_reduce_add)
1596 if (Name.consume_front(
"reduce.")) {
1598 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1599 if (R.match(Name, &
Groups))
1601 .
Case(
"add", Intrinsic::vector_reduce_add)
1602 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1603 .
Case(
"and", Intrinsic::vector_reduce_and)
1604 .
Case(
"or", Intrinsic::vector_reduce_or)
1605 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1606 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1607 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1608 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1609 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1610 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1611 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1616 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1621 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1622 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1627 auto Args =
F->getFunctionType()->params();
1629 {Args[V2 ? 1 : 0]});
1635 if (Name.consume_front(
"splice"))
1639 if (Name.consume_front(
"experimental.stepvector.")) {
1643 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1648 if (Name.starts_with(
"flt.rounds")) {
1651 Intrinsic::get_rounding);
1656 if (Name.starts_with(
"invariant.group.barrier")) {
1658 auto Args =
F->getFunctionType()->params();
1659 Type* ObjectPtr[1] = {Args[0]};
1662 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1667 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1668 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1669 if (IsLifetimeStart || IsLifetimeEnd) {
1670 if (
F->arg_size() == 2) {
1671 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1672 : Intrinsic::lifetime_end;
1677 F->getArg(1)->getType());
1679 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
1699 .StartsWith(
"memcpy.", Intrinsic::memcpy)
1700 .StartsWith(
"memmove.", Intrinsic::memmove)
1702 if (
F->arg_size() == 5) {
1706 F->getFunctionType()->params().slice(0, 3);
1712 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
1715 const auto *FT =
F->getFunctionType();
1716 Type *ParamTypes[2] = {
1717 FT->getParamType(0),
1721 Intrinsic::memset, ParamTypes);
1727 .
StartsWith(
"masked.load", Intrinsic::masked_load)
1728 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
1729 .
StartsWith(
"masked.store", Intrinsic::masked_store)
1730 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
1732 if (MaskedID &&
F->arg_size() == 4) {
1734 if (MaskedID == Intrinsic::masked_load ||
1735 MaskedID == Intrinsic::masked_gather) {
1737 F->getParent(), MaskedID,
1738 {F->getReturnType(), F->getArg(0)->getType()});
1742 F->getParent(), MaskedID,
1743 {F->getArg(0)->getType(), F->getArg(1)->getType()});
1749 if (Name.consume_front(
"nvvm.")) {
1751 if (
F->arg_size() == 1) {
1754 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
1755 .Case(
"clz.i", Intrinsic::ctlz)
1756 .
Case(
"popc.i", Intrinsic::ctpop)
1760 {F->getReturnType()});
1763 }
else if (
F->arg_size() == 2) {
1766 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
1767 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
1768 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
1769 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
1773 {F->getReturnType()});
1779 if (!
F->getReturnType()->getScalarType()->isBFloatTy()) {
1809 F->getParent(), IID,
F->getReturnType(),
1810 F->getFunctionType()->params());
1833 bool Expand =
false;
1834 if (Name.consume_front(
"abs."))
1837 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
1838 else if (Name.consume_front(
"fabs."))
1840 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
1841 else if (Name.consume_front(
"ex2.approx."))
1844 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
1845 else if (Name.consume_front(
"atomic.load."))
1854 else if (Name.consume_front(
"atomic."))
1869 else if (Name.consume_front(
"bitcast."))
1872 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
1873 else if (Name.consume_front(
"rotate."))
1875 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
1876 else if (Name.consume_front(
"ptr.gen.to."))
1879 else if (Name.consume_front(
"ptr."))
1882 else if (Name.consume_front(
"ldg.global."))
1884 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
1885 Name.starts_with(
"p."));
1888 .
Case(
"barrier0",
true)
1889 .
Case(
"barrier.n",
true)
1890 .
Case(
"barrier.sync.cnt",
true)
1891 .
Case(
"barrier.sync",
true)
1892 .
Case(
"barrier",
true)
1893 .
Case(
"bar.sync",
true)
1894 .
Case(
"barrier0.popc",
true)
1895 .
Case(
"barrier0.and",
true)
1896 .
Case(
"barrier0.or",
true)
1897 .
Case(
"clz.ll",
true)
1898 .
Case(
"popc.ll",
true)
1900 .
Case(
"swap.lo.hi.b64",
true)
1901 .
Case(
"tanh.approx.f32",
true)
1913 if (Name.starts_with(
"objectsize.")) {
1914 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
1915 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
1918 Intrinsic::objectsize, Tys);
1925 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
1928 F->getParent(), Intrinsic::ptr_annotation,
1929 {F->arg_begin()->getType(), F->getArg(1)->getType()});
1935 if (Name.consume_front(
"riscv.")) {
1938 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
1939 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
1940 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
1941 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
1944 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
1957 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
1958 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
1967 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
1968 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
1969 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
1970 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
1975 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
1984 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
1986 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
1995 if (Name ==
"stackprotectorcheck") {
2002 if (Name ==
"thread.pointer") {
2004 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2010 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2013 F->getParent(), Intrinsic::var_annotation,
2014 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2017 if (Name.consume_front(
"vector.splice")) {
2018 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2026 if (Name.consume_front(
"wasm.")) {
2029 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2030 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2031 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2036 F->getReturnType());
2040 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2042 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2044 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2063 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2073 std::string
Name =
F->getName().str();
2076 Name,
F->getParent());
2087 if (Result != std::nullopt) {
2103 bool CanUpgradeDebugIntrinsicsToRecords) {
2123 GV->
getName() ==
"llvm.global_dtors")) ||
2138 unsigned N =
Init->getNumOperands();
2139 std::vector<Constant *> NewCtors(
N);
2140 for (
unsigned i = 0; i !=
N; ++i) {
2143 Ctor->getAggregateElement(1),
2157 unsigned NumElts = ResultTy->getNumElements() * 8;
2161 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2171 for (
unsigned l = 0; l != NumElts; l += 16)
2172 for (
unsigned i = 0; i != 16; ++i) {
2173 unsigned Idx = NumElts + i - Shift;
2175 Idx -= NumElts - 16;
2176 Idxs[l + i] = Idx + l;
2179 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2183 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2191 unsigned NumElts = ResultTy->getNumElements() * 8;
2195 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2205 for (
unsigned l = 0; l != NumElts; l += 16)
2206 for (
unsigned i = 0; i != 16; ++i) {
2207 unsigned Idx = i + Shift;
2209 Idx += NumElts - 16;
2210 Idxs[l + i] = Idx + l;
2213 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2217 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2225 Mask = Builder.CreateBitCast(Mask, MaskTy);
2231 for (
unsigned i = 0; i != NumElts; ++i)
2233 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2244 if (
C->isAllOnesValue())
2249 return Builder.CreateSelect(Mask, Op0, Op1);
2256 if (
C->isAllOnesValue())
2260 Mask->getType()->getIntegerBitWidth());
2261 Mask = Builder.CreateBitCast(Mask, MaskTy);
2262 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2263 return Builder.CreateSelect(Mask, Op0, Op1);
2276 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2277 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2282 ShiftVal &= (NumElts - 1);
2291 if (ShiftVal > 16) {
2299 for (
unsigned l = 0; l < NumElts; l += 16) {
2300 for (
unsigned i = 0; i != 16; ++i) {
2301 unsigned Idx = ShiftVal + i;
2302 if (!IsVALIGN && Idx >= 16)
2303 Idx += NumElts - 16;
2304 Indices[l + i] = Idx + l;
2309 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2315 bool ZeroMask,
bool IndexForm) {
2318 unsigned EltWidth = Ty->getScalarSizeInBits();
2319 bool IsFloat = Ty->isFPOrFPVectorTy();
2321 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2322 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2323 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2324 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2325 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2326 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2327 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2328 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2329 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2330 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2331 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2332 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2333 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2334 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2335 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2336 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2337 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2338 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2339 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2340 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2341 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2342 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2343 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2344 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2345 else if (VecWidth == 128 && EltWidth == 16)
2346 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2347 else if (VecWidth == 256 && EltWidth == 16)
2348 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2349 else if (VecWidth == 512 && EltWidth == 16)
2350 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2351 else if (VecWidth == 128 && EltWidth == 8)
2352 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2353 else if (VecWidth == 256 && EltWidth == 8)
2354 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2355 else if (VecWidth == 512 && EltWidth == 8)
2356 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2367 Value *V = Builder.CreateIntrinsic(IID, Args);
2379 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2390 bool IsRotateRight) {
2400 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2401 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2404 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2405 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2450 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2455 bool IsShiftRight,
bool ZeroMask) {
2469 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2470 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2473 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2474 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2489 const Align Alignment =
2491 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2496 if (
C->isAllOnesValue())
2497 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2502 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2508 const Align Alignment =
2517 if (
C->isAllOnesValue())
2518 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2523 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2529 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2530 {Op0, Builder.getInt1(
false)});
2545 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2546 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2547 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2548 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2549 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2552 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2553 LHS = Builder.CreateAnd(
LHS, Mask);
2554 RHS = Builder.CreateAnd(
RHS, Mask);
2571 if (!
C || !
C->isAllOnesValue())
2572 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2577 for (
unsigned i = 0; i != NumElts; ++i)
2579 for (
unsigned i = NumElts; i != 8; ++i)
2580 Indices[i] = NumElts + i % NumElts;
2581 Vec = Builder.CreateShuffleVector(Vec,
2585 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2589 unsigned CC,
bool Signed) {
2597 }
else if (CC == 7) {
2633 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
2634 Value* Cmp = Builder.CreateIsNotNull(AndNode);
2636 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
2637 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
2646 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
2652 Name = Name.substr(12);
2657 if (Name.starts_with(
"max.p")) {
2658 if (VecWidth == 128 && EltWidth == 32)
2659 IID = Intrinsic::x86_sse_max_ps;
2660 else if (VecWidth == 128 && EltWidth == 64)
2661 IID = Intrinsic::x86_sse2_max_pd;
2662 else if (VecWidth == 256 && EltWidth == 32)
2663 IID = Intrinsic::x86_avx_max_ps_256;
2664 else if (VecWidth == 256 && EltWidth == 64)
2665 IID = Intrinsic::x86_avx_max_pd_256;
2668 }
else if (Name.starts_with(
"min.p")) {
2669 if (VecWidth == 128 && EltWidth == 32)
2670 IID = Intrinsic::x86_sse_min_ps;
2671 else if (VecWidth == 128 && EltWidth == 64)
2672 IID = Intrinsic::x86_sse2_min_pd;
2673 else if (VecWidth == 256 && EltWidth == 32)
2674 IID = Intrinsic::x86_avx_min_ps_256;
2675 else if (VecWidth == 256 && EltWidth == 64)
2676 IID = Intrinsic::x86_avx_min_pd_256;
2679 }
else if (Name.starts_with(
"pshuf.b.")) {
2680 if (VecWidth == 128)
2681 IID = Intrinsic::x86_ssse3_pshuf_b_128;
2682 else if (VecWidth == 256)
2683 IID = Intrinsic::x86_avx2_pshuf_b;
2684 else if (VecWidth == 512)
2685 IID = Intrinsic::x86_avx512_pshuf_b_512;
2688 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
2689 if (VecWidth == 128)
2690 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
2691 else if (VecWidth == 256)
2692 IID = Intrinsic::x86_avx2_pmul_hr_sw;
2693 else if (VecWidth == 512)
2694 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
2697 }
else if (Name.starts_with(
"pmulh.w.")) {
2698 if (VecWidth == 128)
2699 IID = Intrinsic::x86_sse2_pmulh_w;
2700 else if (VecWidth == 256)
2701 IID = Intrinsic::x86_avx2_pmulh_w;
2702 else if (VecWidth == 512)
2703 IID = Intrinsic::x86_avx512_pmulh_w_512;
2706 }
else if (Name.starts_with(
"pmulhu.w.")) {
2707 if (VecWidth == 128)
2708 IID = Intrinsic::x86_sse2_pmulhu_w;
2709 else if (VecWidth == 256)
2710 IID = Intrinsic::x86_avx2_pmulhu_w;
2711 else if (VecWidth == 512)
2712 IID = Intrinsic::x86_avx512_pmulhu_w_512;
2715 }
else if (Name.starts_with(
"pmaddw.d.")) {
2716 if (VecWidth == 128)
2717 IID = Intrinsic::x86_sse2_pmadd_wd;
2718 else if (VecWidth == 256)
2719 IID = Intrinsic::x86_avx2_pmadd_wd;
2720 else if (VecWidth == 512)
2721 IID = Intrinsic::x86_avx512_pmaddw_d_512;
2724 }
else if (Name.starts_with(
"pmaddubs.w.")) {
2725 if (VecWidth == 128)
2726 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
2727 else if (VecWidth == 256)
2728 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
2729 else if (VecWidth == 512)
2730 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
2733 }
else if (Name.starts_with(
"packsswb.")) {
2734 if (VecWidth == 128)
2735 IID = Intrinsic::x86_sse2_packsswb_128;
2736 else if (VecWidth == 256)
2737 IID = Intrinsic::x86_avx2_packsswb;
2738 else if (VecWidth == 512)
2739 IID = Intrinsic::x86_avx512_packsswb_512;
2742 }
else if (Name.starts_with(
"packssdw.")) {
2743 if (VecWidth == 128)
2744 IID = Intrinsic::x86_sse2_packssdw_128;
2745 else if (VecWidth == 256)
2746 IID = Intrinsic::x86_avx2_packssdw;
2747 else if (VecWidth == 512)
2748 IID = Intrinsic::x86_avx512_packssdw_512;
2751 }
else if (Name.starts_with(
"packuswb.")) {
2752 if (VecWidth == 128)
2753 IID = Intrinsic::x86_sse2_packuswb_128;
2754 else if (VecWidth == 256)
2755 IID = Intrinsic::x86_avx2_packuswb;
2756 else if (VecWidth == 512)
2757 IID = Intrinsic::x86_avx512_packuswb_512;
2760 }
else if (Name.starts_with(
"packusdw.")) {
2761 if (VecWidth == 128)
2762 IID = Intrinsic::x86_sse41_packusdw;
2763 else if (VecWidth == 256)
2764 IID = Intrinsic::x86_avx2_packusdw;
2765 else if (VecWidth == 512)
2766 IID = Intrinsic::x86_avx512_packusdw_512;
2769 }
else if (Name.starts_with(
"vpermilvar.")) {
2770 if (VecWidth == 128 && EltWidth == 32)
2771 IID = Intrinsic::x86_avx_vpermilvar_ps;
2772 else if (VecWidth == 128 && EltWidth == 64)
2773 IID = Intrinsic::x86_avx_vpermilvar_pd;
2774 else if (VecWidth == 256 && EltWidth == 32)
2775 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
2776 else if (VecWidth == 256 && EltWidth == 64)
2777 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
2778 else if (VecWidth == 512 && EltWidth == 32)
2779 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
2780 else if (VecWidth == 512 && EltWidth == 64)
2781 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
2784 }
else if (Name ==
"cvtpd2dq.256") {
2785 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
2786 }
else if (Name ==
"cvtpd2ps.256") {
2787 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
2788 }
else if (Name ==
"cvttpd2dq.256") {
2789 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
2790 }
else if (Name ==
"cvttps2dq.128") {
2791 IID = Intrinsic::x86_sse2_cvttps2dq;
2792 }
else if (Name ==
"cvttps2dq.256") {
2793 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
2794 }
else if (Name.starts_with(
"permvar.")) {
2796 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2797 IID = Intrinsic::x86_avx2_permps;
2798 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2799 IID = Intrinsic::x86_avx2_permd;
2800 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2801 IID = Intrinsic::x86_avx512_permvar_df_256;
2802 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2803 IID = Intrinsic::x86_avx512_permvar_di_256;
2804 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2805 IID = Intrinsic::x86_avx512_permvar_sf_512;
2806 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2807 IID = Intrinsic::x86_avx512_permvar_si_512;
2808 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2809 IID = Intrinsic::x86_avx512_permvar_df_512;
2810 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2811 IID = Intrinsic::x86_avx512_permvar_di_512;
2812 else if (VecWidth == 128 && EltWidth == 16)
2813 IID = Intrinsic::x86_avx512_permvar_hi_128;
2814 else if (VecWidth == 256 && EltWidth == 16)
2815 IID = Intrinsic::x86_avx512_permvar_hi_256;
2816 else if (VecWidth == 512 && EltWidth == 16)
2817 IID = Intrinsic::x86_avx512_permvar_hi_512;
2818 else if (VecWidth == 128 && EltWidth == 8)
2819 IID = Intrinsic::x86_avx512_permvar_qi_128;
2820 else if (VecWidth == 256 && EltWidth == 8)
2821 IID = Intrinsic::x86_avx512_permvar_qi_256;
2822 else if (VecWidth == 512 && EltWidth == 8)
2823 IID = Intrinsic::x86_avx512_permvar_qi_512;
2826 }
else if (Name.starts_with(
"dbpsadbw.")) {
2827 if (VecWidth == 128)
2828 IID = Intrinsic::x86_avx512_dbpsadbw_128;
2829 else if (VecWidth == 256)
2830 IID = Intrinsic::x86_avx512_dbpsadbw_256;
2831 else if (VecWidth == 512)
2832 IID = Intrinsic::x86_avx512_dbpsadbw_512;
2835 }
else if (Name.starts_with(
"pmultishift.qb.")) {
2836 if (VecWidth == 128)
2837 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
2838 else if (VecWidth == 256)
2839 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
2840 else if (VecWidth == 512)
2841 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
2844 }
else if (Name.starts_with(
"conflict.")) {
2845 if (Name[9] ==
'd' && VecWidth == 128)
2846 IID = Intrinsic::x86_avx512_conflict_d_128;
2847 else if (Name[9] ==
'd' && VecWidth == 256)
2848 IID = Intrinsic::x86_avx512_conflict_d_256;
2849 else if (Name[9] ==
'd' && VecWidth == 512)
2850 IID = Intrinsic::x86_avx512_conflict_d_512;
2851 else if (Name[9] ==
'q' && VecWidth == 128)
2852 IID = Intrinsic::x86_avx512_conflict_q_128;
2853 else if (Name[9] ==
'q' && VecWidth == 256)
2854 IID = Intrinsic::x86_avx512_conflict_q_256;
2855 else if (Name[9] ==
'q' && VecWidth == 512)
2856 IID = Intrinsic::x86_avx512_conflict_q_512;
2859 }
else if (Name.starts_with(
"pavg.")) {
2860 if (Name[5] ==
'b' && VecWidth == 128)
2861 IID = Intrinsic::x86_sse2_pavg_b;
2862 else if (Name[5] ==
'b' && VecWidth == 256)
2863 IID = Intrinsic::x86_avx2_pavg_b;
2864 else if (Name[5] ==
'b' && VecWidth == 512)
2865 IID = Intrinsic::x86_avx512_pavg_b_512;
2866 else if (Name[5] ==
'w' && VecWidth == 128)
2867 IID = Intrinsic::x86_sse2_pavg_w;
2868 else if (Name[5] ==
'w' && VecWidth == 256)
2869 IID = Intrinsic::x86_avx2_pavg_w;
2870 else if (Name[5] ==
'w' && VecWidth == 512)
2871 IID = Intrinsic::x86_avx512_pavg_w_512;
2880 Rep = Builder.CreateIntrinsic(IID, Args);
2891 if (AsmStr->find(
"mov\tfp") == 0 &&
2892 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
2893 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
2894 AsmStr->replace(Pos, 1,
";");
2900 Value *Rep =
nullptr;
2902 if (Name ==
"abs.i" || Name ==
"abs.ll") {
2904 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
2905 {Arg, Builder.getTrue()},
2907 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
2908 Type *Ty = (Name ==
"abs.bf16")
2912 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
2913 Rep = Builder.CreateBitCast(Abs, CI->
getType());
2914 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
2915 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
2916 : Intrinsic::nvvm_fabs;
2917 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
2918 }
else if (Name.consume_front(
"ex2.approx.")) {
2920 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
2921 : Intrinsic::nvvm_ex2_approx;
2922 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
2923 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
2924 Name.starts_with(
"atomic.load.add.f64.p")) {
2927 Rep = Builder.CreateAtomicRMW(
2933 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
2934 Name.starts_with(
"atomic.load.dec.32.p")) {
2939 Rep = Builder.CreateAtomicRMW(
2943 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
2949 Op.contains(
".cta.") ?
"block" :
"");
2950 if (
Op.starts_with(
"cas.")) {
2952 Value *Pair = Builder.CreateAtomicCmpXchg(
2955 Rep = Builder.CreateExtractValue(Pair, 0);
2973 "unexpected nvvm scoped atomic intrinsic");
2974 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
2977 }
else if (Name ==
"clz.ll") {
2980 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
2981 {Arg, Builder.getFalse()},
2983 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
2984 }
else if (Name ==
"popc.ll") {
2988 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
2989 Arg,
nullptr,
"ctpop");
2990 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
2991 }
else if (Name ==
"h2f") {
2993 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
2994 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
2995 }
else if (Name.consume_front(
"bitcast.") &&
2996 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
2999 }
else if (Name ==
"rotate.b32") {
3002 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3003 {Arg, Arg, ShiftAmt});
3004 }
else if (Name ==
"rotate.b64") {
3008 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3009 {Arg, Arg, ZExtShiftAmt});
3010 }
else if (Name ==
"rotate.right.b64") {
3014 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3015 {Arg, Arg, ZExtShiftAmt});
3016 }
else if (Name ==
"swap.lo.hi.b64") {
3019 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3020 {Arg, Arg, Builder.getInt64(32)});
3021 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3024 Name.starts_with(
".to.gen"))) {
3026 }
else if (Name.consume_front(
"ldg.global")) {
3030 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3033 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3035 }
else if (Name ==
"tanh.approx.f32") {
3039 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3041 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3043 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3044 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3046 }
else if (Name ==
"barrier") {
3047 Rep = Builder.CreateIntrinsic(
3048 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3050 }
else if (Name ==
"barrier.sync") {
3051 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3053 }
else if (Name ==
"barrier.sync.cnt") {
3054 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3056 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3057 Name ==
"barrier0.or") {
3059 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3063 .
Case(
"barrier0.popc",
3064 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3065 .
Case(
"barrier0.and",
3066 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3067 .
Case(
"barrier0.or",
3068 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3069 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3070 Rep = Builder.CreateZExt(Bar, CI->
getType());
3074 !
F->getReturnType()->getScalarType()->isBFloatTy()) {
3084 ? Builder.CreateBitCast(Arg, NewType)
3087 Rep = Builder.CreateCall(NewFn, Args);
3088 if (
F->getReturnType()->isIntegerTy())
3089 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3099 Value *Rep =
nullptr;
3101 if (Name.starts_with(
"sse4a.movnt.")) {
3113 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3116 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3117 }
else if (Name.starts_with(
"avx.movnt.") ||
3118 Name.starts_with(
"avx512.storent.")) {
3130 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3131 }
else if (Name ==
"sse2.storel.dq") {
3136 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3137 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3138 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3139 }
else if (Name.starts_with(
"sse.storeu.") ||
3140 Name.starts_with(
"sse2.storeu.") ||
3141 Name.starts_with(
"avx.storeu.")) {
3144 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3145 }
else if (Name ==
"avx512.mask.store.ss") {
3149 }
else if (Name.starts_with(
"avx512.mask.store")) {
3151 bool Aligned = Name[17] !=
'u';
3154 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3157 bool CmpEq = Name[9] ==
'e';
3160 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3161 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3168 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3169 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3171 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3172 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3173 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3174 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3175 Name.starts_with(
"sse2.sqrt.p") ||
3176 Name.starts_with(
"sse.sqrt.p")) {
3177 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3178 {CI->getArgOperand(0)});
3179 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3183 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3184 : Intrinsic::x86_avx512_sqrt_pd_512;
3187 Rep = Builder.CreateIntrinsic(IID, Args);
3189 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3190 {CI->getArgOperand(0)});
3194 }
else if (Name.starts_with(
"avx512.ptestm") ||
3195 Name.starts_with(
"avx512.ptestnm")) {
3199 Rep = Builder.CreateAnd(Op0, Op1);
3205 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3207 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3210 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3213 }
else if (Name.starts_with(
"avx512.kunpck")) {
3218 for (
unsigned i = 0; i != NumElts; ++i)
3227 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3228 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3229 }
else if (Name ==
"avx512.kand.w") {
3232 Rep = Builder.CreateAnd(
LHS,
RHS);
3233 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3234 }
else if (Name ==
"avx512.kandn.w") {
3237 LHS = Builder.CreateNot(
LHS);
3238 Rep = Builder.CreateAnd(
LHS,
RHS);
3239 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3240 }
else if (Name ==
"avx512.kor.w") {
3243 Rep = Builder.CreateOr(
LHS,
RHS);
3244 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3245 }
else if (Name ==
"avx512.kxor.w") {
3248 Rep = Builder.CreateXor(
LHS,
RHS);
3249 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3250 }
else if (Name ==
"avx512.kxnor.w") {
3253 LHS = Builder.CreateNot(
LHS);
3254 Rep = Builder.CreateXor(
LHS,
RHS);
3255 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3256 }
else if (Name ==
"avx512.knot.w") {
3258 Rep = Builder.CreateNot(Rep);
3259 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3260 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3263 Rep = Builder.CreateOr(
LHS,
RHS);
3264 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3266 if (Name[14] ==
'c')
3270 Rep = Builder.CreateICmpEQ(Rep,
C);
3271 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3272 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3273 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3274 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3275 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3278 ConstantInt::get(I32Ty, 0));
3280 ConstantInt::get(I32Ty, 0));
3282 if (Name.contains(
".add."))
3283 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3284 else if (Name.contains(
".sub."))
3285 EltOp = Builder.CreateFSub(Elt0, Elt1);
3286 else if (Name.contains(
".mul."))
3287 EltOp = Builder.CreateFMul(Elt0, Elt1);
3289 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3290 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3291 ConstantInt::get(I32Ty, 0));
3292 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3294 bool CmpEq = Name[16] ==
'e';
3296 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3298 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3305 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3308 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3311 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3318 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3320 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3321 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3323 if (VecWidth == 128 && EltWidth == 32)
3324 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3325 else if (VecWidth == 256 && EltWidth == 32)
3326 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3327 else if (VecWidth == 512 && EltWidth == 32)
3328 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3329 else if (VecWidth == 128 && EltWidth == 64)
3330 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3331 else if (VecWidth == 256 && EltWidth == 64)
3332 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3333 else if (VecWidth == 512 && EltWidth == 64)
3334 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3341 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3344 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3345 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3347 if (VecWidth == 128 && EltWidth == 32)
3348 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3349 else if (VecWidth == 256 && EltWidth == 32)
3350 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3351 else if (VecWidth == 512 && EltWidth == 32)
3352 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3353 else if (VecWidth == 128 && EltWidth == 64)
3354 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3355 else if (VecWidth == 256 && EltWidth == 64)
3356 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3357 else if (VecWidth == 512 && EltWidth == 64)
3358 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3363 if (VecWidth == 512)
3365 Args.push_back(Mask);
3367 Rep = Builder.CreateIntrinsic(IID, Args);
3368 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3372 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3375 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3376 Name.starts_with(
"avx512.cvtw2mask.") ||
3377 Name.starts_with(
"avx512.cvtd2mask.") ||
3378 Name.starts_with(
"avx512.cvtq2mask.")) {
3383 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3384 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3385 Name.starts_with(
"avx512.mask.pabs")) {
3387 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3388 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3389 Name.starts_with(
"avx512.mask.pmaxs")) {
3391 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3392 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3393 Name.starts_with(
"avx512.mask.pmaxu")) {
3395 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3396 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3397 Name.starts_with(
"avx512.mask.pmins")) {
3399 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3400 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3401 Name.starts_with(
"avx512.mask.pminu")) {
3403 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3404 Name ==
"avx512.pmulu.dq.512" ||
3405 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3407 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3408 Name ==
"avx512.pmul.dq.512" ||
3409 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3411 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3412 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3417 }
else if (Name ==
"avx512.cvtusi2sd") {
3422 }
else if (Name ==
"sse2.cvtss2sd") {
3424 Rep = Builder.CreateFPExt(
3427 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3428 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3429 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3430 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3431 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3432 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3433 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3434 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3435 Name ==
"avx512.mask.cvtqq2ps.256" ||
3436 Name ==
"avx512.mask.cvtqq2ps.512" ||
3437 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3438 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3439 Name ==
"avx.cvt.ps2.pd.256" ||
3440 Name ==
"avx512.mask.cvtps2pd.128" ||
3441 Name ==
"avx512.mask.cvtps2pd.256") {
3446 unsigned NumDstElts = DstTy->getNumElements();
3447 if (NumDstElts < SrcTy->getNumElements()) {
3448 assert(NumDstElts == 2 &&
"Unexpected vector size");
3449 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3452 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3453 bool IsUnsigned = Name.contains(
"cvtu");
3455 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3459 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3460 : Intrinsic::x86_avx512_sitofp_round;
3461 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3464 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3465 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3471 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3472 Name.starts_with(
"vcvtph2ps.")) {
3476 unsigned NumDstElts = DstTy->getNumElements();
3477 if (NumDstElts != SrcTy->getNumElements()) {
3478 assert(NumDstElts == 4 &&
"Unexpected vector size");
3479 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3481 Rep = Builder.CreateBitCast(
3483 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3487 }
else if (Name.starts_with(
"avx512.mask.load")) {
3489 bool Aligned = Name[16] !=
'u';
3492 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3496 ResultTy->getNumElements());
3497 Rep = Builder.CreateIntrinsic(
3498 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3500 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3506 Rep = Builder.CreateIntrinsic(
3507 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3509 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3510 Name.starts_with(
"avx512.mask.expand.")) {
3514 ResultTy->getNumElements());
3516 bool IsCompress = Name[12] ==
'c';
3517 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3518 : Intrinsic::x86_avx512_mask_expand;
3519 Rep = Builder.CreateIntrinsic(
3521 }
else if (Name.starts_with(
"xop.vpcom")) {
3523 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3524 Name.ends_with(
"uq"))
3526 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3527 Name.ends_with(
"d") || Name.ends_with(
"q"))
3536 Name = Name.substr(9);
3537 if (Name.starts_with(
"lt"))
3539 else if (Name.starts_with(
"le"))
3541 else if (Name.starts_with(
"gt"))
3543 else if (Name.starts_with(
"ge"))
3545 else if (Name.starts_with(
"eq"))
3547 else if (Name.starts_with(
"ne"))
3549 else if (Name.starts_with(
"false"))
3551 else if (Name.starts_with(
"true"))
3558 }
else if (Name.starts_with(
"xop.vpcmov")) {
3560 Value *NotSel = Builder.CreateNot(Sel);
3563 Rep = Builder.CreateOr(Sel0, Sel1);
3564 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3565 Name.starts_with(
"avx512.mask.prol")) {
3567 }
else if (Name.starts_with(
"avx512.pror") ||
3568 Name.starts_with(
"avx512.mask.pror")) {
3570 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3571 Name.starts_with(
"avx512.mask.vpshld") ||
3572 Name.starts_with(
"avx512.maskz.vpshld")) {
3573 bool ZeroMask = Name[11] ==
'z';
3575 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3576 Name.starts_with(
"avx512.mask.vpshrd") ||
3577 Name.starts_with(
"avx512.maskz.vpshrd")) {
3578 bool ZeroMask = Name[11] ==
'z';
3580 }
else if (Name ==
"sse42.crc32.64.8") {
3583 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3585 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3586 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3587 Name.starts_with(
"avx512.vbroadcast.s")) {
3590 Type *EltTy = VecTy->getElementType();
3591 unsigned EltNum = VecTy->getNumElements();
3595 for (
unsigned I = 0;
I < EltNum; ++
I)
3596 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3597 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3598 Name.starts_with(
"sse41.pmovzx") ||
3599 Name.starts_with(
"avx2.pmovsx") ||
3600 Name.starts_with(
"avx2.pmovzx") ||
3601 Name.starts_with(
"avx512.mask.pmovsx") ||
3602 Name.starts_with(
"avx512.mask.pmovzx")) {
3604 unsigned NumDstElts = DstTy->getNumElements();
3608 for (
unsigned i = 0; i != NumDstElts; ++i)
3613 bool DoSext = Name.contains(
"pmovsx");
3615 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
3620 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
3621 Name ==
"avx512.mask.pmov.qd.512" ||
3622 Name ==
"avx512.mask.pmov.wb.256" ||
3623 Name ==
"avx512.mask.pmov.wb.512") {
3628 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
3629 Name ==
"avx2.vbroadcasti128") {
3635 if (NumSrcElts == 2)
3638 Rep = Builder.CreateShuffleVector(
Load,
3640 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
3641 Name.starts_with(
"avx512.mask.shuf.f")) {
3646 unsigned ControlBitsMask = NumLanes - 1;
3647 unsigned NumControlBits = NumLanes / 2;
3650 for (
unsigned l = 0; l != NumLanes; ++l) {
3651 unsigned LaneMask = (Imm >> (l * NumControlBits)) & ControlBitsMask;
3653 if (l >= NumLanes / 2)
3654 LaneMask += NumLanes;
3655 for (
unsigned i = 0; i != NumElementsInLane; ++i)
3656 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
3662 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
3663 Name.starts_with(
"avx512.mask.broadcasti")) {
3666 unsigned NumDstElts =
3670 for (
unsigned i = 0; i != NumDstElts; ++i)
3671 ShuffleMask[i] = i % NumSrcElts;
3677 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
3678 Name.starts_with(
"avx2.vbroadcast") ||
3679 Name.starts_with(
"avx512.pbroadcast") ||
3680 Name.starts_with(
"avx512.mask.broadcast.s")) {
3687 Rep = Builder.CreateShuffleVector(
Op, M);
3692 }
else if (Name.starts_with(
"sse2.padds.") ||
3693 Name.starts_with(
"avx2.padds.") ||
3694 Name.starts_with(
"avx512.padds.") ||
3695 Name.starts_with(
"avx512.mask.padds.")) {
3697 }
else if (Name.starts_with(
"sse2.psubs.") ||
3698 Name.starts_with(
"avx2.psubs.") ||
3699 Name.starts_with(
"avx512.psubs.") ||
3700 Name.starts_with(
"avx512.mask.psubs.")) {
3702 }
else if (Name.starts_with(
"sse2.paddus.") ||
3703 Name.starts_with(
"avx2.paddus.") ||
3704 Name.starts_with(
"avx512.mask.paddus.")) {
3706 }
else if (Name.starts_with(
"sse2.psubus.") ||
3707 Name.starts_with(
"avx2.psubus.") ||
3708 Name.starts_with(
"avx512.mask.psubus.")) {
3710 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
3715 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
3719 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
3724 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
3729 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
3730 Name ==
"avx512.psll.dq.512") {
3734 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
3735 Name ==
"avx512.psrl.dq.512") {
3739 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
3740 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
3741 Name.starts_with(
"avx2.pblendd.")) {
3746 unsigned NumElts = VecTy->getNumElements();
3749 for (
unsigned i = 0; i != NumElts; ++i)
3750 Idxs[i] = ((Imm >> (i % 8)) & 1) ? i + NumElts : i;
3752 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3753 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
3754 Name ==
"avx2.vinserti128" ||
3755 Name.starts_with(
"avx512.mask.insert")) {
3759 unsigned DstNumElts =
3761 unsigned SrcNumElts =
3763 unsigned Scale = DstNumElts / SrcNumElts;
3770 for (
unsigned i = 0; i != SrcNumElts; ++i)
3772 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
3773 Idxs[i] = SrcNumElts;
3774 Rep = Builder.CreateShuffleVector(Op1, Idxs);
3788 for (
unsigned i = 0; i != DstNumElts; ++i)
3791 for (
unsigned i = 0; i != SrcNumElts; ++i)
3792 Idxs[i + Imm * SrcNumElts] = i + DstNumElts;
3793 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
3799 }
else if (Name.starts_with(
"avx.vextractf128.") ||
3800 Name ==
"avx2.vextracti128" ||
3801 Name.starts_with(
"avx512.mask.vextract")) {
3804 unsigned DstNumElts =
3806 unsigned SrcNumElts =
3808 unsigned Scale = SrcNumElts / DstNumElts;
3815 for (
unsigned i = 0; i != DstNumElts; ++i) {
3816 Idxs[i] = i + (Imm * DstNumElts);
3818 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3824 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
3825 Name.starts_with(
"avx512.mask.perm.di.")) {
3829 unsigned NumElts = VecTy->getNumElements();
3832 for (
unsigned i = 0; i != NumElts; ++i)
3833 Idxs[i] = (i & ~0x3) + ((Imm >> (2 * (i & 0x3))) & 3);
3835 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3840 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
3852 unsigned HalfSize = NumElts / 2;
3864 unsigned StartIndex = (Imm & 0x01) ? HalfSize : 0;
3865 for (
unsigned i = 0; i < HalfSize; ++i)
3866 ShuffleMask[i] = StartIndex + i;
3869 StartIndex = (Imm & 0x10) ? HalfSize : 0;
3870 for (
unsigned i = 0; i < HalfSize; ++i)
3871 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
3873 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
3875 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
3876 Name.starts_with(
"avx512.mask.vpermil.p") ||
3877 Name.starts_with(
"avx512.mask.pshuf.d.")) {
3881 unsigned NumElts = VecTy->getNumElements();
3883 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
3884 unsigned IdxMask = ((1 << IdxSize) - 1);
3890 for (
unsigned i = 0; i != NumElts; ++i)
3891 Idxs[i] = ((Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
3893 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3898 }
else if (Name ==
"sse2.pshufl.w" ||
3899 Name.starts_with(
"avx512.mask.pshufl.w.")) {
3904 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
3908 for (
unsigned l = 0; l != NumElts; l += 8) {
3909 for (
unsigned i = 0; i != 4; ++i)
3910 Idxs[i + l] = ((Imm >> (2 * i)) & 0x3) + l;
3911 for (
unsigned i = 4; i != 8; ++i)
3912 Idxs[i + l] = i + l;
3915 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3920 }
else if (Name ==
"sse2.pshufh.w" ||
3921 Name.starts_with(
"avx512.mask.pshufh.w.")) {
3926 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
3930 for (
unsigned l = 0; l != NumElts; l += 8) {
3931 for (
unsigned i = 0; i != 4; ++i)
3932 Idxs[i + l] = i + l;
3933 for (
unsigned i = 0; i != 4; ++i)
3934 Idxs[i + l + 4] = ((Imm >> (2 * i)) & 0x3) + 4 + l;
3937 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3942 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
3949 unsigned HalfLaneElts = NumLaneElts / 2;
3952 for (
unsigned i = 0; i != NumElts; ++i) {
3954 Idxs[i] = i - (i % NumLaneElts);
3956 if ((i % NumLaneElts) >= HalfLaneElts)
3960 Idxs[i] += (Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
3963 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
3967 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
3968 Name.starts_with(
"avx512.mask.movshdup") ||
3969 Name.starts_with(
"avx512.mask.movsldup")) {
3975 if (Name.starts_with(
"avx512.mask.movshdup."))
3979 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
3980 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
3981 Idxs[i + l + 0] = i + l +
Offset;
3982 Idxs[i + l + 1] = i + l +
Offset;
3985 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
3989 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
3990 Name.starts_with(
"avx512.mask.unpckl.")) {
3997 for (
int l = 0; l != NumElts; l += NumLaneElts)
3998 for (
int i = 0; i != NumLaneElts; ++i)
3999 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4001 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4005 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4006 Name.starts_with(
"avx512.mask.unpckh.")) {
4013 for (
int l = 0; l != NumElts; l += NumLaneElts)
4014 for (
int i = 0; i != NumLaneElts; ++i)
4015 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4017 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4021 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4022 Name.starts_with(
"avx512.mask.pand.")) {
4025 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4027 Rep = Builder.CreateBitCast(Rep, FTy);
4030 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4031 Name.starts_with(
"avx512.mask.pandn.")) {
4034 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4035 Rep = Builder.CreateAnd(Rep,
4037 Rep = Builder.CreateBitCast(Rep, FTy);
4040 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4041 Name.starts_with(
"avx512.mask.por.")) {
4044 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4046 Rep = Builder.CreateBitCast(Rep, FTy);
4049 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4050 Name.starts_with(
"avx512.mask.pxor.")) {
4053 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4055 Rep = Builder.CreateBitCast(Rep, FTy);
4058 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4062 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4066 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4070 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4071 if (Name.ends_with(
".512")) {
4073 if (Name[17] ==
's')
4074 IID = Intrinsic::x86_avx512_add_ps_512;
4076 IID = Intrinsic::x86_avx512_add_pd_512;
4078 Rep = Builder.CreateIntrinsic(
4086 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4087 if (Name.ends_with(
".512")) {
4089 if (Name[17] ==
's')
4090 IID = Intrinsic::x86_avx512_div_ps_512;
4092 IID = Intrinsic::x86_avx512_div_pd_512;
4094 Rep = Builder.CreateIntrinsic(
4102 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4103 if (Name.ends_with(
".512")) {
4105 if (Name[17] ==
's')
4106 IID = Intrinsic::x86_avx512_mul_ps_512;
4108 IID = Intrinsic::x86_avx512_mul_pd_512;
4110 Rep = Builder.CreateIntrinsic(
4118 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4119 if (Name.ends_with(
".512")) {
4121 if (Name[17] ==
's')
4122 IID = Intrinsic::x86_avx512_sub_ps_512;
4124 IID = Intrinsic::x86_avx512_sub_pd_512;
4126 Rep = Builder.CreateIntrinsic(
4134 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4135 Name.starts_with(
"avx512.mask.min.p")) &&
4136 Name.drop_front(18) ==
".512") {
4137 bool IsDouble = Name[17] ==
'd';
4138 bool IsMin = Name[13] ==
'i';
4140 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4141 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4144 Rep = Builder.CreateIntrinsic(
4149 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4151 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4152 {CI->getArgOperand(0), Builder.getInt1(false)});
4155 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4156 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4157 bool IsVariable = Name[16] ==
'v';
4158 char Size = Name[16] ==
'.' ? Name[17]
4159 : Name[17] ==
'.' ? Name[18]
4160 : Name[18] ==
'.' ? Name[19]
4164 if (IsVariable && Name[17] !=
'.') {
4165 if (
Size ==
'd' && Name[17] ==
'2')
4166 IID = Intrinsic::x86_avx2_psllv_q;
4167 else if (
Size ==
'd' && Name[17] ==
'4')
4168 IID = Intrinsic::x86_avx2_psllv_q_256;
4169 else if (
Size ==
's' && Name[17] ==
'4')
4170 IID = Intrinsic::x86_avx2_psllv_d;
4171 else if (
Size ==
's' && Name[17] ==
'8')
4172 IID = Intrinsic::x86_avx2_psllv_d_256;
4173 else if (
Size ==
'h' && Name[17] ==
'8')
4174 IID = Intrinsic::x86_avx512_psllv_w_128;
4175 else if (
Size ==
'h' && Name[17] ==
'1')
4176 IID = Intrinsic::x86_avx512_psllv_w_256;
4177 else if (Name[17] ==
'3' && Name[18] ==
'2')
4178 IID = Intrinsic::x86_avx512_psllv_w_512;
4181 }
else if (Name.ends_with(
".128")) {
4183 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4184 : Intrinsic::x86_sse2_psll_d;
4185 else if (
Size ==
'q')
4186 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4187 : Intrinsic::x86_sse2_psll_q;
4188 else if (
Size ==
'w')
4189 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4190 : Intrinsic::x86_sse2_psll_w;
4193 }
else if (Name.ends_with(
".256")) {
4195 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4196 : Intrinsic::x86_avx2_psll_d;
4197 else if (
Size ==
'q')
4198 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4199 : Intrinsic::x86_avx2_psll_q;
4200 else if (
Size ==
'w')
4201 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4202 : Intrinsic::x86_avx2_psll_w;
4207 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4208 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4209 : Intrinsic::x86_avx512_psll_d_512;
4210 else if (
Size ==
'q')
4211 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4212 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4213 : Intrinsic::x86_avx512_psll_q_512;
4214 else if (
Size ==
'w')
4215 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4216 : Intrinsic::x86_avx512_psll_w_512;
4222 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4223 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4224 bool IsVariable = Name[16] ==
'v';
4225 char Size = Name[16] ==
'.' ? Name[17]
4226 : Name[17] ==
'.' ? Name[18]
4227 : Name[18] ==
'.' ? Name[19]
4231 if (IsVariable && Name[17] !=
'.') {
4232 if (
Size ==
'd' && Name[17] ==
'2')
4233 IID = Intrinsic::x86_avx2_psrlv_q;
4234 else if (
Size ==
'd' && Name[17] ==
'4')
4235 IID = Intrinsic::x86_avx2_psrlv_q_256;
4236 else if (
Size ==
's' && Name[17] ==
'4')
4237 IID = Intrinsic::x86_avx2_psrlv_d;
4238 else if (
Size ==
's' && Name[17] ==
'8')
4239 IID = Intrinsic::x86_avx2_psrlv_d_256;
4240 else if (
Size ==
'h' && Name[17] ==
'8')
4241 IID = Intrinsic::x86_avx512_psrlv_w_128;
4242 else if (
Size ==
'h' && Name[17] ==
'1')
4243 IID = Intrinsic::x86_avx512_psrlv_w_256;
4244 else if (Name[17] ==
'3' && Name[18] ==
'2')
4245 IID = Intrinsic::x86_avx512_psrlv_w_512;
4248 }
else if (Name.ends_with(
".128")) {
4250 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4251 : Intrinsic::x86_sse2_psrl_d;
4252 else if (
Size ==
'q')
4253 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4254 : Intrinsic::x86_sse2_psrl_q;
4255 else if (
Size ==
'w')
4256 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4257 : Intrinsic::x86_sse2_psrl_w;
4260 }
else if (Name.ends_with(
".256")) {
4262 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4263 : Intrinsic::x86_avx2_psrl_d;
4264 else if (
Size ==
'q')
4265 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4266 : Intrinsic::x86_avx2_psrl_q;
4267 else if (
Size ==
'w')
4268 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4269 : Intrinsic::x86_avx2_psrl_w;
4274 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4275 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4276 : Intrinsic::x86_avx512_psrl_d_512;
4277 else if (
Size ==
'q')
4278 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4279 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4280 : Intrinsic::x86_avx512_psrl_q_512;
4281 else if (
Size ==
'w')
4282 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4283 : Intrinsic::x86_avx512_psrl_w_512;
4289 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4290 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4291 bool IsVariable = Name[16] ==
'v';
4292 char Size = Name[16] ==
'.' ? Name[17]
4293 : Name[17] ==
'.' ? Name[18]
4294 : Name[18] ==
'.' ? Name[19]
4298 if (IsVariable && Name[17] !=
'.') {
4299 if (
Size ==
's' && Name[17] ==
'4')
4300 IID = Intrinsic::x86_avx2_psrav_d;
4301 else if (
Size ==
's' && Name[17] ==
'8')
4302 IID = Intrinsic::x86_avx2_psrav_d_256;
4303 else if (
Size ==
'h' && Name[17] ==
'8')
4304 IID = Intrinsic::x86_avx512_psrav_w_128;
4305 else if (
Size ==
'h' && Name[17] ==
'1')
4306 IID = Intrinsic::x86_avx512_psrav_w_256;
4307 else if (Name[17] ==
'3' && Name[18] ==
'2')
4308 IID = Intrinsic::x86_avx512_psrav_w_512;
4311 }
else if (Name.ends_with(
".128")) {
4313 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4314 : Intrinsic::x86_sse2_psra_d;
4315 else if (
Size ==
'q')
4316 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4317 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4318 : Intrinsic::x86_avx512_psra_q_128;
4319 else if (
Size ==
'w')
4320 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4321 : Intrinsic::x86_sse2_psra_w;
4324 }
else if (Name.ends_with(
".256")) {
4326 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4327 : Intrinsic::x86_avx2_psra_d;
4328 else if (
Size ==
'q')
4329 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4330 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4331 : Intrinsic::x86_avx512_psra_q_256;
4332 else if (
Size ==
'w')
4333 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4334 : Intrinsic::x86_avx2_psra_w;
4339 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4340 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4341 : Intrinsic::x86_avx512_psra_d_512;
4342 else if (
Size ==
'q')
4343 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4344 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4345 : Intrinsic::x86_avx512_psra_q_512;
4346 else if (
Size ==
'w')
4347 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4348 : Intrinsic::x86_avx512_psra_w_512;
4354 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4356 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4358 }
else if (Name.ends_with(
".movntdqa")) {
4362 LoadInst *LI = Builder.CreateAlignedLoad(
4367 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4368 Name.starts_with(
"fma.vfmsub.") ||
4369 Name.starts_with(
"fma.vfnmadd.") ||
4370 Name.starts_with(
"fma.vfnmsub.")) {
4371 bool NegMul = Name[6] ==
'n';
4372 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4373 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4384 if (NegMul && !IsScalar)
4385 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4386 if (NegMul && IsScalar)
4387 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4389 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4391 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4395 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4403 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4407 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4408 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4409 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4410 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4411 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4412 bool IsMask3 = Name[11] ==
'3';
4413 bool IsMaskZ = Name[11] ==
'z';
4415 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4416 bool NegMul = Name[2] ==
'n';
4417 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4423 if (NegMul && (IsMask3 || IsMaskZ))
4424 A = Builder.CreateFNeg(
A);
4425 if (NegMul && !(IsMask3 || IsMaskZ))
4426 B = Builder.CreateFNeg(
B);
4428 C = Builder.CreateFNeg(
C);
4430 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4431 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4432 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4439 if (Name.back() ==
'd')
4440 IID = Intrinsic::x86_avx512_vfmadd_f64;
4442 IID = Intrinsic::x86_avx512_vfmadd_f32;
4443 Rep = Builder.CreateIntrinsic(IID,
Ops);
4445 Rep = Builder.CreateFMA(
A,
B,
C);
4454 if (NegAcc && IsMask3)
4459 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4461 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4462 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4463 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4464 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4465 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4466 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4467 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4468 bool IsMask3 = Name[11] ==
'3';
4469 bool IsMaskZ = Name[11] ==
'z';
4471 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4472 bool NegMul = Name[2] ==
'n';
4473 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4479 if (NegMul && (IsMask3 || IsMaskZ))
4480 A = Builder.CreateFNeg(
A);
4481 if (NegMul && !(IsMask3 || IsMaskZ))
4482 B = Builder.CreateFNeg(
B);
4484 C = Builder.CreateFNeg(
C);
4491 if (Name[Name.size() - 5] ==
's')
4492 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4494 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4498 Rep = Builder.CreateFMA(
A,
B,
C);
4506 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4510 if (VecWidth == 128 && EltWidth == 32)
4511 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4512 else if (VecWidth == 256 && EltWidth == 32)
4513 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4514 else if (VecWidth == 128 && EltWidth == 64)
4515 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4516 else if (VecWidth == 256 && EltWidth == 64)
4517 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4523 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4524 Rep = Builder.CreateIntrinsic(IID,
Ops);
4525 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4526 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4527 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4528 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4529 bool IsMask3 = Name[11] ==
'3';
4530 bool IsMaskZ = Name[11] ==
'z';
4532 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4533 bool IsSubAdd = Name[3] ==
's';
4537 if (Name[Name.size() - 5] ==
's')
4538 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4540 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4545 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4547 Rep = Builder.CreateIntrinsic(IID,
Ops);
4556 Value *Odd = Builder.CreateCall(FMA,
Ops);
4557 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4558 Value *Even = Builder.CreateCall(FMA,
Ops);
4564 for (
int i = 0; i != NumElts; ++i)
4565 Idxs[i] = i + (i % 2) * NumElts;
4567 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4575 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4576 Name.starts_with(
"avx512.maskz.pternlog.")) {
4577 bool ZeroMask = Name[11] ==
'z';
4581 if (VecWidth == 128 && EltWidth == 32)
4582 IID = Intrinsic::x86_avx512_pternlog_d_128;
4583 else if (VecWidth == 256 && EltWidth == 32)
4584 IID = Intrinsic::x86_avx512_pternlog_d_256;
4585 else if (VecWidth == 512 && EltWidth == 32)
4586 IID = Intrinsic::x86_avx512_pternlog_d_512;
4587 else if (VecWidth == 128 && EltWidth == 64)
4588 IID = Intrinsic::x86_avx512_pternlog_q_128;
4589 else if (VecWidth == 256 && EltWidth == 64)
4590 IID = Intrinsic::x86_avx512_pternlog_q_256;
4591 else if (VecWidth == 512 && EltWidth == 64)
4592 IID = Intrinsic::x86_avx512_pternlog_q_512;
4598 Rep = Builder.CreateIntrinsic(IID, Args);
4602 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
4603 Name.starts_with(
"avx512.maskz.vpmadd52")) {
4604 bool ZeroMask = Name[11] ==
'z';
4605 bool High = Name[20] ==
'h' || Name[21] ==
'h';
4608 if (VecWidth == 128 && !
High)
4609 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
4610 else if (VecWidth == 256 && !
High)
4611 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
4612 else if (VecWidth == 512 && !
High)
4613 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
4614 else if (VecWidth == 128 &&
High)
4615 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
4616 else if (VecWidth == 256 &&
High)
4617 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
4618 else if (VecWidth == 512 &&
High)
4619 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
4625 Rep = Builder.CreateIntrinsic(IID, Args);
4629 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
4630 Name.starts_with(
"avx512.mask.vpermt2var.") ||
4631 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
4632 bool ZeroMask = Name[11] ==
'z';
4633 bool IndexForm = Name[17] ==
'i';
4635 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
4636 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
4637 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
4638 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
4639 bool ZeroMask = Name[11] ==
'z';
4640 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4643 if (VecWidth == 128 && !IsSaturating)
4644 IID = Intrinsic::x86_avx512_vpdpbusd_128;
4645 else if (VecWidth == 256 && !IsSaturating)
4646 IID = Intrinsic::x86_avx512_vpdpbusd_256;
4647 else if (VecWidth == 512 && !IsSaturating)
4648 IID = Intrinsic::x86_avx512_vpdpbusd_512;
4649 else if (VecWidth == 128 && IsSaturating)
4650 IID = Intrinsic::x86_avx512_vpdpbusds_128;
4651 else if (VecWidth == 256 && IsSaturating)
4652 IID = Intrinsic::x86_avx512_vpdpbusds_256;
4653 else if (VecWidth == 512 && IsSaturating)
4654 IID = Intrinsic::x86_avx512_vpdpbusds_512;
4664 if (Args[1]->
getType()->isVectorTy() &&
4667 ->isIntegerTy(32) &&
4668 Args[2]->
getType()->isVectorTy() &&
4671 ->isIntegerTy(32)) {
4672 Type *NewArgType =
nullptr;
4673 if (VecWidth == 128)
4675 else if (VecWidth == 256)
4677 else if (VecWidth == 512)
4683 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4684 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4687 Rep = Builder.CreateIntrinsic(IID, Args);
4691 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
4692 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
4693 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
4694 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
4695 bool ZeroMask = Name[11] ==
'z';
4696 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
4699 if (VecWidth == 128 && !IsSaturating)
4700 IID = Intrinsic::x86_avx512_vpdpwssd_128;
4701 else if (VecWidth == 256 && !IsSaturating)
4702 IID = Intrinsic::x86_avx512_vpdpwssd_256;
4703 else if (VecWidth == 512 && !IsSaturating)
4704 IID = Intrinsic::x86_avx512_vpdpwssd_512;
4705 else if (VecWidth == 128 && IsSaturating)
4706 IID = Intrinsic::x86_avx512_vpdpwssds_128;
4707 else if (VecWidth == 256 && IsSaturating)
4708 IID = Intrinsic::x86_avx512_vpdpwssds_256;
4709 else if (VecWidth == 512 && IsSaturating)
4710 IID = Intrinsic::x86_avx512_vpdpwssds_512;
4720 if (Args[1]->
getType()->isVectorTy() &&
4723 ->isIntegerTy(32) &&
4724 Args[2]->
getType()->isVectorTy() &&
4727 ->isIntegerTy(32)) {
4728 Type *NewArgType =
nullptr;
4729 if (VecWidth == 128)
4731 else if (VecWidth == 256)
4733 else if (VecWidth == 512)
4739 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
4740 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
4743 Rep = Builder.CreateIntrinsic(IID, Args);
4747 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
4748 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
4749 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
4751 if (Name[0] ==
'a' && Name.back() ==
'2')
4752 IID = Intrinsic::x86_addcarry_32;
4753 else if (Name[0] ==
'a' && Name.back() ==
'4')
4754 IID = Intrinsic::x86_addcarry_64;
4755 else if (Name[0] ==
's' && Name.back() ==
'2')
4756 IID = Intrinsic::x86_subborrow_32;
4757 else if (Name[0] ==
's' && Name.back() ==
'4')
4758 IID = Intrinsic::x86_subborrow_64;
4765 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
4768 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
4771 Value *CF = Builder.CreateExtractValue(NewCall, 0);
4775 }
else if (Name.starts_with(
"avx512.mask.") &&
4778 }
else if (Name.starts_with(
"bmi.pdep.")) {
4780 }
else if (Name.starts_with(
"bmi.pext.")) {
4790 if (Name.starts_with(
"neon.bfcvt")) {
4791 if (Name.starts_with(
"neon.bfcvtn2")) {
4793 std::iota(LoMask.
begin(), LoMask.
end(), 0);
4795 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4796 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
4799 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
4800 }
else if (Name.starts_with(
"neon.bfcvtn")) {
4802 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
4806 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
4807 return Builder.CreateShuffleVector(
4810 return Builder.CreateFPTrunc(CI->
getOperand(0),
4813 }
else if (Name.starts_with(
"sve.fcvt")) {
4816 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
4817 .
Case(
"sve.fcvtnt.bf16f32",
4818 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
4830 if (Args[1]->
getType() != BadPredTy)
4833 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
4834 BadPredTy, Args[1]);
4835 Args[1] = Builder.CreateIntrinsic(
4836 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
4838 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
4842 if (Name ==
"neon.vcvtfp2hf")
4843 return Builder.CreateBitCast(
4844 Builder.CreateFPTrunc(
4848 if (Name ==
"neon.vcvthf2fp")
4849 return Builder.CreateFPExt(
4850 Builder.CreateBitCast(
4860 if (Name ==
"mve.vctp64.old") {
4863 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
4866 Value *C1 = Builder.CreateIntrinsic(
4867 Intrinsic::arm_mve_pred_v2i,
4869 return Builder.CreateIntrinsic(
4870 Intrinsic::arm_mve_pred_i2v,
4872 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
4873 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
4874 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
4875 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
4877 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
4878 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
4879 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
4880 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
4882 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
4883 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
4884 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
4885 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
4886 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
4887 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
4888 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
4889 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
4890 std::vector<Type *> Tys;
4894 case Intrinsic::arm_mve_mull_int_predicated:
4895 case Intrinsic::arm_mve_vqdmull_predicated:
4896 case Intrinsic::arm_mve_vldr_gather_base_predicated:
4899 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
4900 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
4901 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
4905 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
4909 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
4913 case Intrinsic::arm_cde_vcx1q_predicated:
4914 case Intrinsic::arm_cde_vcx1qa_predicated:
4915 case Intrinsic::arm_cde_vcx2q_predicated:
4916 case Intrinsic::arm_cde_vcx2qa_predicated:
4917 case Intrinsic::arm_cde_vcx3q_predicated:
4918 case Intrinsic::arm_cde_vcx3qa_predicated:
4925 std::vector<Value *>
Ops;
4927 Type *Ty =
Op->getType();
4928 if (Ty->getScalarSizeInBits() == 1) {
4929 Value *C1 = Builder.CreateIntrinsic(
4930 Intrinsic::arm_mve_pred_v2i,
4932 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
4937 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
4952 auto UpgradeLegacyWMMAIUIntrinsicCall =
4957 Args.push_back(Builder.getFalse());
4961 F->getParent(),
F->getIntrinsicID(), OverloadTys);
4968 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
4973 NewCall->copyMetadata(*CI);
4977 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
4978 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
4979 "intrinsic should have 7 arguments");
4982 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
4984 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
4985 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
4986 "intrinsic should have 8 arguments");
4991 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
4994 switch (
F->getIntrinsicID()) {
4997 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
4998 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
4999 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5000 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5001 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5002 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5017 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5020 F->getParent(),
F->getIntrinsicID(), Overloads);
5025 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5030 NewCall->copyMetadata(*CI);
5031 NewCall->takeName(CI);
5053 if (NumOperands < 3)
5066 bool IsVolatile =
false;
5070 if (NumOperands > 3)
5075 if (NumOperands > 5) {
5077 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5091 if (VT->getElementType()->isIntegerTy(16)) {
5094 Val = Builder.CreateBitCast(Val, AsBF16);
5102 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5104 unsigned AddrSpace = PtrTy->getAddressSpace();
5107 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5109 RMW->
setMetadata(
"amdgpu.ignore.denormal.mode", EmptyMD);
5114 MDNode *RangeNotPrivate =
5117 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5123 return Builder.CreateBitCast(RMW, RetTy);
5144 return MAV->getMetadata();
5153 if (Name ==
"label") {
5155 }
else if (Name ==
"assign") {
5162 }
else if (Name ==
"declare") {
5166 }
else if (Name ==
"addr") {
5176 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5177 }
else if (Name ==
"value") {
5180 unsigned ExprOp = 2;
5195 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5203 int64_t OffsetVal =
Offset->getSExtValue();
5204 return Builder.CreateIntrinsic(OffsetVal >= 0
5205 ? Intrinsic::vector_splice_left
5206 : Intrinsic::vector_splice_right,
5208 {CI->getArgOperand(0), CI->getArgOperand(1),
5209 Builder.getInt32(std::abs(OffsetVal))});
5214 if (Name.starts_with(
"to.fp16")) {
5216 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5217 return Builder.CreateBitCast(Cast, CI->
getType());
5220 if (Name.starts_with(
"from.fp16")) {
5222 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5223 return Builder.CreateFPExt(Cast, CI->
getType());
5234 if (Defaults.empty())
5237 unsigned OldArgCount = CI->
arg_size();
5238 unsigned NewArgCount = NewFn->
arg_size();
5242 if (OldArgCount >= NewArgCount)
5250 if (OldArgCount < FirstDefault)
5255 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5256 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5257 "missing argument outside the default range");
5258 Type *ParamTy = NewFT->getParamType(Idx);
5263 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5269 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5301 if (!Name.consume_front(
"llvm."))
5304 bool IsX86 = Name.consume_front(
"x86.");
5305 bool IsNVVM = Name.consume_front(
"nvvm.");
5306 bool IsAArch64 = Name.consume_front(
"aarch64.");
5307 bool IsARM = Name.consume_front(
"arm.");
5308 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5309 bool IsDbg = Name.consume_front(
"dbg.");
5311 (Name.consume_front(
"experimental.vector.splice") ||
5312 Name.consume_front(
"vector.splice")) &&
5313 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5314 Value *Rep =
nullptr;
5316 if (!IsX86 && Name ==
"stackprotectorcheck") {
5318 }
else if (IsNVVM) {
5322 }
else if (IsAArch64) {
5326 }
else if (IsAMDGCN) {
5330 }
else if (IsOldSplice) {
5332 }
else if (Name.consume_front(
"convert.")) {
5334 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5347 const auto &DefaultCase = [&]() ->
void {
5355 "Unknown function for CallBase upgrade and isn't just a name change");
5363 "Return type must have changed");
5364 assert(OldST->getNumElements() ==
5366 "Must have same number of elements");
5369 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5372 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5373 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5374 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5398 case Intrinsic::arm_neon_vst1:
5399 case Intrinsic::arm_neon_vst2:
5400 case Intrinsic::arm_neon_vst3:
5401 case Intrinsic::arm_neon_vst4:
5402 case Intrinsic::arm_neon_vst2lane:
5403 case Intrinsic::arm_neon_vst3lane:
5404 case Intrinsic::arm_neon_vst4lane: {
5406 NewCall = Builder.CreateCall(NewFn, Args);
5409 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5410 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5411 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5416 NewCall = Builder.CreateCall(NewFn, Args);
5419 case Intrinsic::aarch64_sve_ld3_sret:
5420 case Intrinsic::aarch64_sve_ld4_sret:
5421 case Intrinsic::aarch64_sve_ld2_sret: {
5429 Name = Name.substr(5);
5436 unsigned MinElts = RetTy->getMinNumElements() /
N;
5438 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5440 for (
unsigned I = 0;
I <
N;
I++) {
5441 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5442 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5448 case Intrinsic::coro_end_async:
5449 case Intrinsic::coro_end: {
5451 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5453 NewCall = Builder.CreateCall(NewFn, Args);
5458 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5459 Value *InRamp = Builder.CreateCall(IsInRamp);
5469 case Intrinsic::vector_extract: {
5471 Name = Name.substr(5);
5472 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5477 unsigned MinElts = RetTy->getMinNumElements();
5480 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5484 case Intrinsic::vector_insert: {
5486 Name = Name.substr(5);
5487 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5491 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5496 NewCall = Builder.CreateCall(
5500 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
5506 assert(
N > 1 &&
"Create is expected to be between 2-4");
5509 unsigned MinElts = RetTy->getMinNumElements() /
N;
5510 for (
unsigned I = 0;
I <
N;
I++) {
5512 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
5519 case Intrinsic::arm_neon_bfdot:
5520 case Intrinsic::arm_neon_bfmmla:
5521 case Intrinsic::arm_neon_bfmlalb:
5522 case Intrinsic::arm_neon_bfmlalt:
5523 case Intrinsic::aarch64_neon_bfdot:
5524 case Intrinsic::aarch64_neon_bfmmla:
5525 case Intrinsic::aarch64_neon_bfmlalb:
5526 case Intrinsic::aarch64_neon_bfmlalt: {
5529 "Mismatch between function args and call args");
5530 size_t OperandWidth =
5532 assert((OperandWidth == 64 || OperandWidth == 128) &&
5533 "Unexpected operand width");
5535 auto Iter = CI->
args().begin();
5536 Args.push_back(*Iter++);
5537 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5538 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
5539 NewCall = Builder.CreateCall(NewFn, Args);
5543 case Intrinsic::bitreverse:
5544 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5547 case Intrinsic::ctlz:
5548 case Intrinsic::cttz: {
5555 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
5559 case Intrinsic::objectsize: {
5560 Value *NullIsUnknownSize =
5564 NewCall = Builder.CreateCall(
5569 case Intrinsic::ctpop:
5570 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
5572 case Intrinsic::dbg_value: {
5574 Name = Name.substr(5);
5576 if (Name.starts_with(
"dbg.addr")) {
5590 if (
Offset->isNullValue()) {
5591 NewCall = Builder.CreateCall(
5600 case Intrinsic::ptr_annotation:
5608 NewCall = Builder.CreateCall(
5617 case Intrinsic::var_annotation:
5624 NewCall = Builder.CreateCall(
5633 case Intrinsic::riscv_aes32dsi:
5634 case Intrinsic::riscv_aes32dsmi:
5635 case Intrinsic::riscv_aes32esi:
5636 case Intrinsic::riscv_aes32esmi:
5637 case Intrinsic::riscv_sm4ks:
5638 case Intrinsic::riscv_sm4ed: {
5648 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
5649 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
5655 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
5656 Value *Res = NewCall;
5658 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5664 case Intrinsic::nvvm_mapa_shared_cluster: {
5668 Value *Res = NewCall;
5669 Res = Builder.CreateAddrSpaceCast(
5676 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster:
5677 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
5680 Args[0] = Builder.CreateAddrSpaceCast(
5683 NewCall = Builder.CreateCall(NewFn, Args);
5689 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d:
5690 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d:
5691 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d:
5692 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d:
5693 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d:
5694 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d:
5695 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d:
5696 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d: {
5703 Args[0] = Builder.CreateAddrSpaceCast(
5712 Args.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
5714 NewCall = Builder.CreateCall(NewFn, Args);
5720 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
5721 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
5722 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
5723 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
5724 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
5725 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
5726 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
5727 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
5729 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
5733 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
5734 NewCall = Builder.CreateCall(NewFn, Args);
5737 case Intrinsic::nvvm_tcgen05_mma_shared:
5738 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5739 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5740 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
5741 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
5742 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
5743 case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
5744 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
5745 case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
5746 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
5747 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
5748 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
5749 case Intrinsic::nvvm_tcgen05_mma_sp_shared:
5750 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
5751 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
5752 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
5753 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
5754 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
5755 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
5756 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
5757 case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
5758 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
5759 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
5760 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
5761 case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
5762 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
5763 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
5764 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
5765 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
5766 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
5767 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
5768 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
5769 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
5770 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
5771 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
5772 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
5773 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
5774 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
5775 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
5777 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
5778 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
5780 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
5781 case Intrinsic::nvvm_tcgen05_mma_tensor:
5782 case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
5783 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
5784 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
5785 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
5786 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
5787 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
5788 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
5789 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
5790 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
5791 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
5792 case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
5793 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
5794 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
5795 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
5797 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
5798 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
5800 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
5802 Args.push_back(Builder.getInt32(0));
5803 NewCall = Builder.CreateCall(NewFn, Args);
5806 case Intrinsic::riscv_sha256sig0:
5807 case Intrinsic::riscv_sha256sig1:
5808 case Intrinsic::riscv_sha256sum0:
5809 case Intrinsic::riscv_sha256sum1:
5810 case Intrinsic::riscv_sm3p0:
5811 case Intrinsic::riscv_sm3p1: {
5818 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
5820 NewCall = Builder.CreateCall(NewFn, Arg);
5822 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
5829 case Intrinsic::x86_xop_vfrcz_ss:
5830 case Intrinsic::x86_xop_vfrcz_sd:
5831 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
5834 case Intrinsic::x86_xop_vpermil2pd:
5835 case Intrinsic::x86_xop_vpermil2ps:
5836 case Intrinsic::x86_xop_vpermil2pd_256:
5837 case Intrinsic::x86_xop_vpermil2ps_256: {
5841 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
5842 NewCall = Builder.CreateCall(NewFn, Args);
5846 case Intrinsic::x86_sse41_ptestc:
5847 case Intrinsic::x86_sse41_ptestz:
5848 case Intrinsic::x86_sse41_ptestnzc: {
5862 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
5863 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
5865 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
5869 case Intrinsic::x86_rdtscp: {
5875 NewCall = Builder.CreateCall(NewFn);
5877 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5880 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
5888 case Intrinsic::x86_sse41_insertps:
5889 case Intrinsic::x86_sse41_dppd:
5890 case Intrinsic::x86_sse41_dpps:
5891 case Intrinsic::x86_sse41_mpsadbw:
5892 case Intrinsic::x86_avx_dp_ps_256:
5893 case Intrinsic::x86_avx2_mpsadbw: {
5899 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
5900 NewCall = Builder.CreateCall(NewFn, Args);
5904 case Intrinsic::x86_avx512_mask_cmp_pd_128:
5905 case Intrinsic::x86_avx512_mask_cmp_pd_256:
5906 case Intrinsic::x86_avx512_mask_cmp_pd_512:
5907 case Intrinsic::x86_avx512_mask_cmp_ps_128:
5908 case Intrinsic::x86_avx512_mask_cmp_ps_256:
5909 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
5915 NewCall = Builder.CreateCall(NewFn, Args);
5924 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
5925 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
5926 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
5927 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
5928 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
5929 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
5933 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
5934 Args[1] = Builder.CreateBitCast(
5937 NewCall = Builder.CreateCall(NewFn, Args);
5938 Value *Res = Builder.CreateBitCast(
5946 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
5947 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
5948 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
5952 Args[1] = Builder.CreateBitCast(
5954 Args[2] = Builder.CreateBitCast(
5957 NewCall = Builder.CreateCall(NewFn, Args);
5961 case Intrinsic::thread_pointer: {
5962 NewCall = Builder.CreateCall(NewFn, {});
5966 case Intrinsic::memcpy:
5967 case Intrinsic::memmove:
5968 case Intrinsic::memset: {
5984 NewCall = Builder.CreateCall(NewFn, Args);
5986 AttributeList NewAttrs = AttributeList::get(
5987 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
5988 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
5989 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
5994 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
5997 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6001 case Intrinsic::masked_load:
6002 case Intrinsic::masked_gather:
6003 case Intrinsic::masked_store:
6004 case Intrinsic::masked_scatter: {
6010 auto GetMaybeAlign = [](
Value *
Op) {
6020 auto GetAlign = [&](
Value *
Op) {
6029 case Intrinsic::masked_load:
6030 NewCall = Builder.CreateMaskedLoad(
6034 case Intrinsic::masked_gather:
6035 NewCall = Builder.CreateMaskedGather(
6041 case Intrinsic::masked_store:
6042 NewCall = Builder.CreateMaskedStore(
6046 case Intrinsic::masked_scatter:
6047 NewCall = Builder.CreateMaskedScatter(
6049 DL.getValueOrABITypeAlignment(
6063 case Intrinsic::lifetime_start:
6064 case Intrinsic::lifetime_end: {
6076 NewCall = Builder.CreateLifetimeStart(Ptr);
6078 NewCall = Builder.CreateLifetimeEnd(Ptr);
6087 case Intrinsic::x86_avx512_vpdpbusd_128:
6088 case Intrinsic::x86_avx512_vpdpbusd_256:
6089 case Intrinsic::x86_avx512_vpdpbusd_512:
6090 case Intrinsic::x86_avx512_vpdpbusds_128:
6091 case Intrinsic::x86_avx512_vpdpbusds_256:
6092 case Intrinsic::x86_avx512_vpdpbusds_512:
6093 case Intrinsic::x86_avx2_vpdpbssd_128:
6094 case Intrinsic::x86_avx2_vpdpbssd_256:
6095 case Intrinsic::x86_avx10_vpdpbssd_512:
6096 case Intrinsic::x86_avx2_vpdpbssds_128:
6097 case Intrinsic::x86_avx2_vpdpbssds_256:
6098 case Intrinsic::x86_avx10_vpdpbssds_512:
6099 case Intrinsic::x86_avx2_vpdpbsud_128:
6100 case Intrinsic::x86_avx2_vpdpbsud_256:
6101 case Intrinsic::x86_avx10_vpdpbsud_512:
6102 case Intrinsic::x86_avx2_vpdpbsuds_128:
6103 case Intrinsic::x86_avx2_vpdpbsuds_256:
6104 case Intrinsic::x86_avx10_vpdpbsuds_512:
6105 case Intrinsic::x86_avx2_vpdpbuud_128:
6106 case Intrinsic::x86_avx2_vpdpbuud_256:
6107 case Intrinsic::x86_avx10_vpdpbuud_512:
6108 case Intrinsic::x86_avx2_vpdpbuuds_128:
6109 case Intrinsic::x86_avx2_vpdpbuuds_256:
6110 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6115 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6116 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6118 NewCall = Builder.CreateCall(NewFn, Args);
6121 case Intrinsic::x86_avx512_vpdpwssd_128:
6122 case Intrinsic::x86_avx512_vpdpwssd_256:
6123 case Intrinsic::x86_avx512_vpdpwssd_512:
6124 case Intrinsic::x86_avx512_vpdpwssds_128:
6125 case Intrinsic::x86_avx512_vpdpwssds_256:
6126 case Intrinsic::x86_avx512_vpdpwssds_512:
6127 case Intrinsic::x86_avx2_vpdpwsud_128:
6128 case Intrinsic::x86_avx2_vpdpwsud_256:
6129 case Intrinsic::x86_avx10_vpdpwsud_512:
6130 case Intrinsic::x86_avx2_vpdpwsuds_128:
6131 case Intrinsic::x86_avx2_vpdpwsuds_256:
6132 case Intrinsic::x86_avx10_vpdpwsuds_512:
6133 case Intrinsic::x86_avx2_vpdpwusd_128:
6134 case Intrinsic::x86_avx2_vpdpwusd_256:
6135 case Intrinsic::x86_avx10_vpdpwusd_512:
6136 case Intrinsic::x86_avx2_vpdpwusds_128:
6137 case Intrinsic::x86_avx2_vpdpwusds_256:
6138 case Intrinsic::x86_avx10_vpdpwusds_512:
6139 case Intrinsic::x86_avx2_vpdpwuud_128:
6140 case Intrinsic::x86_avx2_vpdpwuud_256:
6141 case Intrinsic::x86_avx10_vpdpwuud_512:
6142 case Intrinsic::x86_avx2_vpdpwuuds_128:
6143 case Intrinsic::x86_avx2_vpdpwuuds_256:
6144 case Intrinsic::x86_avx10_vpdpwuuds_512:
6149 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6150 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6152 NewCall = Builder.CreateCall(NewFn, Args);
6155 assert(NewCall &&
"Should have either set this variable or returned through "
6156 "the default case");
6163 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6177 F->eraseFromParent();
6183 if (NumOperands == 0)
6191 if (NumOperands == 3) {
6195 Metadata *Elts2[] = {ScalarType, ScalarType,
6209 if (
Opc != Instruction::BitCast)
6213 Type *SrcTy = V->getType();
6230 if (
Opc != Instruction::BitCast)
6233 Type *SrcTy =
C->getType();
6250 if (Flag.getNumOperands() < 3)
6251 return std::nullopt;
6253 return Name->getString();
6254 return std::nullopt;
6268 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6269 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6270 if (auto Name = getModuleFlagNameSafely(*Flag))
6271 return *Name ==
"Debug Info Version";
6274 if (OpIt != ModFlags->op_end()) {
6275 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6282 bool BrokenDebugInfo =
false;
6285 if (!BrokenDebugInfo)
6291 M.getContext().diagnose(Diag);
6298 M.getContext().diagnose(DiagVersion);
6308 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6311 if (
F->hasFnAttribute(Attr)) {
6314 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6316 auto [Part, Rest] = S.
split(
',');
6322 const unsigned Dim = DimC -
'x';
6323 assert(Dim < 3 &&
"Unexpected dim char");
6333 F->addFnAttr(Attr, NewAttr);
6337 return S ==
"x" || S ==
"y" || S ==
"z";
6342 if (K ==
"kernel") {
6354 const unsigned Idx = (AlignIdxValuePair >> 16);
6355 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6360 if (K ==
"maxclusterrank" || K ==
"cluster_max_blocks") {
6365 if (K ==
"minctasm") {
6370 if (K ==
"maxnreg") {
6375 if (K.consume_front(
"maxntid") &&
isXYZ(K)) {
6379 if (K.consume_front(
"reqntid") &&
isXYZ(K)) {
6383 if (K.consume_front(
"cluster_dim_") &&
isXYZ(K)) {
6387 if (K ==
"grid_constant") {
6402 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6409 if (!SeenNodes.
insert(MD).second)
6416 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6423 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6425 const MDOperand &V = MD->getOperand(j + 1);
6428 NewOperands.
append({K, V});
6431 if (NewOperands.
size() > 1)
6444 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6445 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6446 if (ModRetainReleaseMarker) {
6452 ID->getString().split(ValueComp,
"#");
6453 if (ValueComp.
size() == 2) {
6454 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6458 M.eraseNamedMetadata(ModRetainReleaseMarker);
6469 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
6495 bool InvalidCast =
false;
6497 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
6510 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
6512 Args.push_back(Arg);
6519 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
6524 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
6537 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
6545 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
6546 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
6547 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
6548 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
6549 {
"objc_autoreleaseReturnValue",
6550 llvm::Intrinsic::objc_autoreleaseReturnValue},
6551 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
6552 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
6553 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
6554 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
6555 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
6556 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
6557 {
"objc_release", llvm::Intrinsic::objc_release},
6558 {
"objc_retain", llvm::Intrinsic::objc_retain},
6559 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
6560 {
"objc_retainAutoreleaseReturnValue",
6561 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
6562 {
"objc_retainAutoreleasedReturnValue",
6563 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
6564 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
6565 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
6566 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
6567 {
"objc_unsafeClaimAutoreleasedReturnValue",
6568 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
6569 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
6570 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
6571 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
6572 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
6573 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
6574 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
6575 {
"objc_arc_annotation_topdown_bbstart",
6576 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
6577 {
"objc_arc_annotation_topdown_bbend",
6578 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
6579 {
"objc_arc_annotation_bottomup_bbstart",
6580 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
6581 {
"objc_arc_annotation_bottomup_bbend",
6582 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
6584 for (
auto &
I : RuntimeFuncs)
6585 UpgradeToIntrinsic(
I.first,
I.second);
6609 std::optional<bool> UseAddressDisc;
6612 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6613 for (
const MDNode *Flag : ModFlags->operands()) {
6615 if (Name && (*Name ==
"ptrauth-init-fini" ||
6616 *Name ==
"ptrauth-init-fini-address-discrimination"))
6621 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
6622 constexpr unsigned ExpectedConstDisc = 0xD9D4;
6623 constexpr unsigned ExpectedAddressMarker = 1;
6626 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
6629 bool HasAddressDisc;
6630 if (!CPA->hasAddressDiscriminator())
6631 HasAddressDisc =
false;
6632 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
6633 HasAddressDisc =
true;
6637 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
6640 UseAddressDisc = HasAddressDisc;
6641 return CPA->getPointer();
6645 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
6648 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
6650 if (!GV || !GV->hasInitializer())
6654 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
6657 std::vector<Constant *> NewStructors;
6658 NewStructors.reserve(OldStructorsArray->getNumOperands());
6660 for (
Use &U : OldStructorsArray->operands()) {
6669 Func = UpgradeSinglePointer(Func);
6673 NewStructors.push_back(
6682 if (GlobalArraysToUpgrade.
empty())
6684 assert(UseAddressDisc.has_value());
6686 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
6687 GV->setInitializer(NewInit);
6690 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
6700 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
6704 bool HasObjCFlag =
false, HasClassProperties =
false;
6705 bool HasSwiftVersionFlag =
false;
6706 uint8_t SwiftMajorVersion, SwiftMinorVersion;
6713 if (
Op->getNumOperands() != 3)
6727 if (ID->getString() ==
"Objective-C Image Info Version")
6729 if (ID->getString() ==
"Objective-C Class Properties")
6730 HasClassProperties =
true;
6732 if (ID->getString() ==
"PIC Level") {
6733 if (
auto *Behavior =
6735 uint64_t V = Behavior->getLimitedValue();
6741 if (ID->getString() ==
"PIE Level")
6742 if (
auto *Behavior =
6749 if (ID->getString() ==
"branch-target-enforcement" ||
6750 ID->getString().starts_with(
"sign-return-address")) {
6751 if (
auto *Behavior =
6757 Op->getOperand(1),
Op->getOperand(2)};
6767 if (ID->getString() ==
"Objective-C Image Info Section") {
6770 Value->getString().split(ValueComp,
" ");
6771 if (ValueComp.
size() != 1) {
6772 std::string NewValue;
6773 for (
auto &S : ValueComp)
6774 NewValue += S.str();
6785 if (ID->getString() ==
"Objective-C Garbage Collection") {
6788 assert(Md->getValue() &&
"Expected non-empty metadata");
6789 auto Type = Md->getValue()->getType();
6792 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
6793 if ((Val & 0xff) != Val) {
6794 HasSwiftVersionFlag =
true;
6795 SwiftABIVersion = (Val & 0xff00) >> 8;
6796 SwiftMajorVersion = (Val & 0xff000000) >> 24;
6797 SwiftMinorVersion = (Val & 0xff0000) >> 16;
6808 if (ID->getString() ==
"amdgpu_code_object_version") {
6811 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
6820 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
6849 if (HasObjCFlag && !HasClassProperties) {
6855 if (HasSwiftVersionFlag) {
6859 ConstantInt::get(Int8Ty, SwiftMajorVersion));
6861 ConstantInt::get(Int8Ty, SwiftMinorVersion));
6869 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
6873 auto MatchesVersion = [](
const MDNode *
Op) {
6874 return Op->getNumOperands() >= 3 &&
6888 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
6889 assert(
Op->getNumOperands() >= 2 &&
6890 "Expected at least 2 operands - name and linkage type");
6902 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
6913 auto TrimSpaces = [](
StringRef Section) -> std::string {
6915 Section.split(Components,
',');
6920 for (
auto Component : Components)
6921 OS <<
',' << Component.trim();
6926 for (
auto &GV : M.globals()) {
6927 if (!GV.hasSection())
6932 if (!Section.starts_with(
"__DATA, __objc_catlist"))
6937 GV.setSection(TrimSpaces(Section));
6953struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
6954 StrictFPUpgradeVisitor() =
default;
6957 if (!
Call.isStrictFP())
6963 Call.removeFnAttr(Attribute::StrictFP);
6964 Call.addFnAttr(Attribute::NoBuiltin);
6969struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
6970 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
6971 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
6973 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
6988 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
6989 StrictFPUpgradeVisitor SFPV;
6994 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
6995 F.getReturnType(),
F.getAttributes().getRetAttrs()));
6996 for (
auto &Arg :
F.args())
6998 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7000 bool AddingAttrs =
false, RemovingAttrs =
false;
7001 AttrBuilder AttrsToAdd(
F.getContext());
7006 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7007 A.isValid() &&
A.isStringAttribute()) {
7008 F.setSection(
A.getValueAsString());
7010 RemovingAttrs =
true;
7014 A.isValid() &&
A.isStringAttribute()) {
7017 AddingAttrs = RemovingAttrs =
true;
7020 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7021 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7023 RemovingAttrs =
true;
7024 if (
A.getValueAsString() ==
"true") {
7025 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7034 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7037 if (
A.getValueAsBool()) {
7038 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7044 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7045 RemovingAttrs =
true;
7052 bool HandleDenormalMode =
false;
7054 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7057 DenormalFPMath = ParsedMode;
7059 AddingAttrs = RemovingAttrs =
true;
7060 HandleDenormalMode =
true;
7064 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7068 DenormalFPMathF32 = ParsedMode;
7070 AddingAttrs = RemovingAttrs =
true;
7071 HandleDenormalMode =
true;
7075 if (HandleDenormalMode)
7076 AttrsToAdd.addDenormalFPEnvAttr(
7080 F.removeFnAttrs(AttrsToRemove);
7083 F.addFnAttrs(AttrsToAdd);
7089 if (!
F.hasFnAttribute(FnAttrName))
7090 F.addFnAttr(FnAttrName,
Value);
7097 if (!
F.hasFnAttribute(FnAttrName)) {
7099 F.addFnAttr(FnAttrName);
7101 auto A =
F.getFnAttribute(FnAttrName);
7102 if (
"false" ==
A.getValueAsString())
7103 F.removeFnAttr(FnAttrName);
7104 else if (
"true" ==
A.getValueAsString()) {
7105 F.removeFnAttr(FnAttrName);
7106 F.addFnAttr(FnAttrName);
7112 Triple T(M.getTargetTriple());
7113 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7123 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7127 if (
Op->getNumOperands() != 3)
7136 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7137 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7138 : IDStr ==
"guarded-control-stack" ? &GCSValue
7139 : IDStr ==
"sign-return-address" ? &SRAValue
7140 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7141 : IDStr ==
"sign-return-address-with-bkey"
7147 *ValPtr = CI->getZExtValue();
7153 bool BTE = BTEValue == 1;
7154 bool BPPLR = BPPLRValue == 1;
7155 bool GCS = GCSValue == 1;
7156 bool SRA = SRAValue == 1;
7159 if (SRA && SRAALLValue == 1)
7160 SignTypeValue =
"all";
7163 if (SRA && SRABKeyValue == 1)
7164 SignKeyValue =
"b_key";
7166 for (
Function &
F : M.getFunctionList()) {
7167 if (
F.isDeclaration())
7174 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7175 A.isValid() &&
"none" ==
A.getValueAsString()) {
7176 F.removeFnAttr(
"sign-return-address");
7177 F.removeFnAttr(
"sign-return-address-key");
7193 if (SRAALLValue == 1)
7195 if (SRABKeyValue == 1)
7222 if (
T->getNumOperands() < 1)
7227 if (S->getString().starts_with(
"llvm.vectorizer."))
7233 StringRef OldPrefix =
"llvm.vectorizer.";
7236 if (OldTag ==
"llvm.vectorizer.unroll")
7248 if (
T->getNumOperands() < 1)
7260 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7273 Ops.reserve(
T->getNumOperands());
7274 Ops.push_back(NewTag);
7275 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7276 Ops.push_back(
T->getOperand(
I));
7293 if (
T->isDistinct()) {
7294 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7306 Ops.reserve(
T->getNumOperands());
7317 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7318 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7319 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7322 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7324 auto I =
DL.find(
"-n64-");
7326 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7331 std::string Res =
DL.str();
7334 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7335 Res.append(Res.empty() ?
"G1" :
"-G1");
7343 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7344 Res.append(
"-ni:7:8:9");
7346 if (
DL.ends_with(
"ni:7"))
7348 if (
DL.ends_with(
"ni:7:8"))
7353 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7354 Res.append(
"-p7:160:256:256:32");
7355 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7356 Res.append(
"-p8:128:128:128:48");
7357 constexpr StringRef OldP8(
"-p8:128:128-");
7358 if (
DL.contains(OldP8))
7359 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7360 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7361 Res.append(
"-p9:192:256:256:32");
7365 if (!
DL.contains(
"m:e"))
7366 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7371 if (
T.isSystemZ() && !
DL.empty()) {
7373 if (!
DL.contains(
"-S64"))
7374 return "E-S64" +
DL.drop_front(1).str();
7378 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7381 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7382 if (!
DL.contains(AddrSpaces)) {
7384 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7385 if (R.match(Res, &
Groups))
7391 if (
T.isAArch64()) {
7393 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7394 Res.append(
"-Fn32");
7395 AddPtr32Ptr64AddrSpaces();
7399 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7403 std::string I64 =
"-i64:64";
7404 std::string I128 =
"-i128:128";
7406 size_t Pos = Res.find(I64);
7407 if (Pos !=
size_t(-1))
7408 Res.insert(Pos + I64.size(), I128);
7412 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7413 size_t Pos = Res.find(
"-S128");
7416 Res.insert(Pos,
"-f64:32:64");
7422 AddPtr32Ptr64AddrSpaces();
7430 if (!
T.isOSIAMCU()) {
7431 std::string I128 =
"-i128:128";
7434 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7435 if (R.match(Res, &
Groups))
7443 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
7445 auto I =
Ref.find(
"-f80:32-");
7447 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
7455 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
7458 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
7459 B.removeAttribute(
"no-frame-pointer-elim");
7461 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
7463 if (FramePointer !=
"all")
7464 FramePointer =
"non-leaf";
7465 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
7467 if (!FramePointer.
empty())
7468 B.addAttribute(
"frame-pointer", FramePointer);
7470 A =
B.getAttribute(
"null-pointer-is-valid");
7473 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
7474 B.removeAttribute(
"null-pointer-is-valid");
7475 if (NullPointerIsValid)
7476 B.addAttribute(Attribute::NullPointerIsValid);
7479 A =
B.getAttribute(
"uniform-work-group-size");
7483 bool IsTrue = Val ==
"true";
7484 B.removeAttribute(
"uniform-work-group-size");
7486 B.addAttribute(
"uniform-work-group-size");
7497 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static Metadata * upgradeLoopArgument(Metadata *MD)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static 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 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 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 Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05MMAIntrinsic(Function *F, StringRef Name)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static cl::opt< bool > DisableAutoUpgradeDebugInfo("disable-auto-upgrade-debug-info", cl::desc("Disable autoupgrade of debug info"))
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static 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.
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ ICMP_ULT
unsigned less than
@ ICMP_SGE
signed greater or equal
@ ICMP_ULE
unsigned less or equal
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Type * getValueType() const
const Constant * getInitializer() const
getInitializer - Return the initializer for this global variable.
bool hasInitializer() const
Definitions have initializers, declarations don't.
PointerType * getPtrTy(unsigned AddrSpace=0)
Fetch the type representing a pointer.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Base class for instruction visitors.
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
LLVM_ABI const Module * getModule() const
Return the module owning the function this instruction belongs to or nullptr it the function does not...
LLVM_ABI InstListType::iterator eraseFromParent()
This method unlinks 'this' from the containing basic block and deletes it.
LLVM_ABI void setMetadata(unsigned KindID, MDNode *Node)
Set the metadata of the specified kind to the specified node.
LLVM_ABI FastMathFlags getFastMathFlags() const LLVM_READONLY
Convenience function for getting all the fast-math flags, which must be an operator which supports th...
LLVM_ABI void copyMetadata(const Instruction &SrcInst, ArrayRef< unsigned > WL=ArrayRef< unsigned >())
Copy metadata from SrcInst to this instruction.
LLVM_ABI const DataLayout & getDataLayout() const
Get the data layout of the module this instruction belongs to.
This is an important class for using LLVM in a threaded context.
LLVM_ABI SyncScope::ID getOrInsertSyncScopeID(StringRef SSN)
getOrInsertSyncScopeID - Maps synchronization scope name to synchronization scope ID.
An instruction for reading from memory.
LLVM_ABI MDNode * createRange(const APInt &Lo, const APInt &Hi)
Return metadata describing the range [Lo, Hi).
const MDOperand & getOperand(unsigned I) const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
A switch()-like statement whose cases are string literals.
StringSwitch & Case(StringLiteral S, T Value)
StringSwitch & StartsWith(StringLiteral S, T Value)
StringSwitch & Cases(std::initializer_list< StringLiteral > CaseStrings, T Value)
Class to represent struct types.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
unsigned getNumElements() const
Random access to the elements.
Type * getElementType(unsigned N) const
The TimeTraceScope is a helper class to call the begin and end functions of the time trace profiler.
Triple - Helper class for working with autoconf configuration names.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
static LLVM_ABI IntegerType * getInt64Ty(LLVMContext &C)
bool isVectorTy() const
True if this is an instance of VectorType.
static LLVM_ABI IntegerType * getInt32Ty(LLVMContext &C)
bool isFloatTy() const
Return true if this is 'float', a 32-bit IEEE fp type.
bool isBFloatTy() const
Return true if this is 'bfloat', a 16-bit bfloat type.
LLVM_ABI unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
static LLVM_ABI IntegerType * getInt8Ty(LLVMContext &C)
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
LLVM_ABI TypeSize getPrimitiveSizeInBits() const LLVM_READONLY
Return the basic size of this type if it is a primitive type.
static LLVM_ABI IntegerType * getInt16Ty(LLVMContext &C)
LLVM_ABI unsigned getScalarSizeInBits() const LLVM_READONLY
If this is a vector type, return the getPrimitiveSizeInBits value for the element type.
bool isPtrOrPtrVectorTy() const
Return true if this is a pointer type or a vector of pointer types.
bool isIntegerTy() const
True if this is an instance of IntegerType.
bool isFPOrFPVectorTy() const
Return true if this is a FP type or a vector of FP.
static LLVM_ABI Type * getFloatTy(LLVMContext &C)
static LLVM_ABI Type * getBFloatTy(LLVMContext &C)
static LLVM_ABI Type * getHalfTy(LLVMContext &C)
bool isVoidTy() const
Return true if this is 'void'.
A Use represents the edge between a Value definition and its users.
Value * getOperand(unsigned i) const
unsigned getNumOperands() const
LLVM Value Representation.
Type * getType() const
All values are typed, get the type of this value.
LLVM_ABI void print(raw_ostream &O, bool IsForDebug=false) const
Implement operator<< on Value.
LLVM_ABI void setName(const Twine &Name)
Change the name of the value.
LLVM_ABI void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
LLVMContext & getContext() const
All values hold a context through their type.
iterator_range< user_iterator > users()
LLVM_ABI const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
LLVM_ABI void takeName(Value *V)
Transfer the name from V to this value.
Base class of all SIMD vector types.
static VectorType * getInteger(VectorType *VTy)
This static method gets a VectorType with the same number of elements as the input type,...
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
constexpr ScalarTy getFixedValue() const
const ParentTy * getParent() const
self_iterator getIterator()
A raw_ostream that writes to an SmallVector or SmallString.
StringRef str() const
Return a StringRef for the vector contents.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ LOCAL_ADDRESS
Address space for local memory.
@ FLAT_ADDRESS
Address space for flat memory.
@ PRIVATE_ADDRESS
Address space for private memory.
@ PTX_Kernel
Call to a PTX kernel. Passes all arguments in parameter space.
std::optional< ABIType > parseABIType(StringRef S)
Parse the string spelling used by the "float-abi" IR module flag into an ABIType.
LLVM_ABI std::optional< Function * > remangleIntrinsicFunction(Function *F)
LLVM_ABI Function * getOrInsertDeclaration(Module *M, ID id, ArrayRef< Type * > OverloadTys={})
Look up the Function declaration of the intrinsic id in the Module M.
LLVM_ABI ID lookupIntrinsicID(StringRef Name)
This does the actual lookup of an intrinsic ID which matches the given function name.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI 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.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
@ Dynamic
Denotes mode unknown at compile time.
ArrayRef(const T &OneElt) -> ArrayRef< T >
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
@ Default
The result value is uniform if and only if all operands are uniform.
LLVM_ABI bool verifyModule(const Module &M, raw_ostream *OS=nullptr, bool *BrokenDebugInfo=nullptr)
Check a module for errors.
LLVM_ABI void reportFatalUsageError(Error Err)
Report a fatal error that does not indicate a bug in LLVM.
void swap(llvm::BitVector &LHS, llvm::BitVector &RHS)
Implement std::swap in terms of BitVector swap.
This struct is a compact representation of a valid (non-zero power of two) alignment.
Represents the full denormal controls for a function, including the default mode and the f32 specific...
Represent subnormal handling kind for floating point instruction inputs and outputs.
static constexpr DenormalMode getInvalid()
constexpr bool isValid() const
static constexpr DenormalMode getIEEE()
This struct is a compact representation of a valid (power of two) or undefined (0) alignment.