36#include "llvm/IR/IntrinsicsAArch64.h"
37#include "llvm/IR/IntrinsicsAMDGPU.h"
38#include "llvm/IR/IntrinsicsARM.h"
39#include "llvm/IR/IntrinsicsNVPTX.h"
40#include "llvm/IR/IntrinsicsRISCV.h"
41#include "llvm/IR/IntrinsicsWebAssembly.h"
42#include "llvm/IR/IntrinsicsX86.h"
65#define DEBUG_TYPE "auto-upgrade"
84 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
99 Type *LastArgType =
F->getFunctionType()->getParamType(
100 F->getFunctionType()->getNumParams() - 1);
115 if (
F->getReturnType()->isVectorTy())
128 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
129 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
146 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
147 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
161 if (
F->getReturnType()->getScalarType()->isBFloatTy())
171 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
185 if (Name.consume_front(
"avx."))
186 return (Name.starts_with(
"blend.p") ||
187 Name ==
"cvt.ps2.pd.256" ||
188 Name ==
"cvtdq2.pd.256" ||
189 Name ==
"cvtdq2.ps.256" ||
190 Name.starts_with(
"movnt.") ||
191 Name.starts_with(
"sqrt.p") ||
192 Name.starts_with(
"storeu.") ||
193 Name.starts_with(
"vbroadcast.s") ||
194 Name.starts_with(
"vbroadcastf128") ||
195 Name.starts_with(
"vextractf128.") ||
196 Name.starts_with(
"vinsertf128.") ||
197 Name.starts_with(
"vperm2f128.") ||
198 Name.starts_with(
"vpermil."));
200 if (Name.consume_front(
"avx2."))
201 return (Name ==
"movntdqa" ||
202 Name.starts_with(
"pabs.") ||
203 Name.starts_with(
"padds.") ||
204 Name.starts_with(
"paddus.") ||
205 Name.starts_with(
"pblendd.") ||
207 Name.starts_with(
"pbroadcast") ||
208 Name.starts_with(
"pcmpeq.") ||
209 Name.starts_with(
"pcmpgt.") ||
210 Name.starts_with(
"pmax") ||
211 Name.starts_with(
"pmin") ||
212 Name.starts_with(
"pmovsx") ||
213 Name.starts_with(
"pmovzx") ||
214 Name.starts_with(
"pmulh.w") ||
215 Name.starts_with(
"pmulhu.w") ||
217 Name ==
"pmulu.dq" ||
218 Name.starts_with(
"psll.dq") ||
219 Name.starts_with(
"psrl.dq") ||
220 Name.starts_with(
"psubs.") ||
221 Name.starts_with(
"psubus.") ||
222 Name.starts_with(
"vbroadcast") ||
223 Name ==
"vbroadcasti128" ||
224 Name ==
"vextracti128" ||
225 Name ==
"vinserti128" ||
226 Name ==
"vperm2i128");
228 if (Name.consume_front(
"avx512.")) {
229 if (Name.consume_front(
"mask."))
231 return (Name.starts_with(
"add.p") ||
232 Name.starts_with(
"and.") ||
233 Name.starts_with(
"andn.") ||
234 Name.starts_with(
"broadcast.s") ||
235 Name.starts_with(
"broadcastf32x4.") ||
236 Name.starts_with(
"broadcastf32x8.") ||
237 Name.starts_with(
"broadcastf64x2.") ||
238 Name.starts_with(
"broadcastf64x4.") ||
239 Name.starts_with(
"broadcasti32x4.") ||
240 Name.starts_with(
"broadcasti32x8.") ||
241 Name.starts_with(
"broadcasti64x2.") ||
242 Name.starts_with(
"broadcasti64x4.") ||
243 Name.starts_with(
"cmp.b") ||
244 Name.starts_with(
"cmp.d") ||
245 Name.starts_with(
"cmp.q") ||
246 Name.starts_with(
"cmp.w") ||
247 Name.starts_with(
"compress.b") ||
248 Name.starts_with(
"compress.d") ||
249 Name.starts_with(
"compress.p") ||
250 Name.starts_with(
"compress.q") ||
251 Name.starts_with(
"compress.store.") ||
252 Name.starts_with(
"compress.w") ||
253 Name.starts_with(
"conflict.") ||
254 Name.starts_with(
"cvtdq2pd.") ||
255 Name.starts_with(
"cvtdq2ps.") ||
256 Name ==
"cvtpd2dq.256" ||
257 Name ==
"cvtpd2ps.256" ||
258 Name ==
"cvtps2pd.128" ||
259 Name ==
"cvtps2pd.256" ||
260 Name.starts_with(
"cvtqq2pd.") ||
261 Name ==
"cvtqq2ps.256" ||
262 Name ==
"cvtqq2ps.512" ||
263 Name ==
"cvttpd2dq.256" ||
264 Name ==
"cvttps2dq.128" ||
265 Name ==
"cvttps2dq.256" ||
266 Name.starts_with(
"cvtudq2pd.") ||
267 Name.starts_with(
"cvtudq2ps.") ||
268 Name.starts_with(
"cvtuqq2pd.") ||
269 Name ==
"cvtuqq2ps.256" ||
270 Name ==
"cvtuqq2ps.512" ||
271 Name.starts_with(
"dbpsadbw.") ||
272 Name.starts_with(
"div.p") ||
273 Name.starts_with(
"expand.b") ||
274 Name.starts_with(
"expand.d") ||
275 Name.starts_with(
"expand.load.") ||
276 Name.starts_with(
"expand.p") ||
277 Name.starts_with(
"expand.q") ||
278 Name.starts_with(
"expand.w") ||
279 Name.starts_with(
"fpclass.p") ||
280 Name.starts_with(
"insert") ||
281 Name.starts_with(
"load.") ||
282 Name.starts_with(
"loadu.") ||
283 Name.starts_with(
"lzcnt.") ||
284 Name.starts_with(
"max.p") ||
285 Name.starts_with(
"min.p") ||
286 Name.starts_with(
"movddup") ||
287 Name.starts_with(
"move.s") ||
288 Name.starts_with(
"movshdup") ||
289 Name.starts_with(
"movsldup") ||
290 Name.starts_with(
"mul.p") ||
291 Name.starts_with(
"or.") ||
292 Name.starts_with(
"pabs.") ||
293 Name.starts_with(
"packssdw.") ||
294 Name.starts_with(
"packsswb.") ||
295 Name.starts_with(
"packusdw.") ||
296 Name.starts_with(
"packuswb.") ||
297 Name.starts_with(
"padd.") ||
298 Name.starts_with(
"padds.") ||
299 Name.starts_with(
"paddus.") ||
300 Name.starts_with(
"palignr.") ||
301 Name.starts_with(
"pand.") ||
302 Name.starts_with(
"pandn.") ||
303 Name.starts_with(
"pavg") ||
304 Name.starts_with(
"pbroadcast") ||
305 Name.starts_with(
"pcmpeq.") ||
306 Name.starts_with(
"pcmpgt.") ||
307 Name.starts_with(
"perm.df.") ||
308 Name.starts_with(
"perm.di.") ||
309 Name.starts_with(
"permvar.") ||
310 Name.starts_with(
"pmaddubs.w.") ||
311 Name.starts_with(
"pmaddw.d.") ||
312 Name.starts_with(
"pmax") ||
313 Name.starts_with(
"pmin") ||
314 Name ==
"pmov.qd.256" ||
315 Name ==
"pmov.qd.512" ||
316 Name ==
"pmov.wb.256" ||
317 Name ==
"pmov.wb.512" ||
318 Name.starts_with(
"pmovsx") ||
319 Name.starts_with(
"pmovzx") ||
320 Name.starts_with(
"pmul.dq.") ||
321 Name.starts_with(
"pmul.hr.sw.") ||
322 Name.starts_with(
"pmulh.w.") ||
323 Name.starts_with(
"pmulhu.w.") ||
324 Name.starts_with(
"pmull.") ||
325 Name.starts_with(
"pmultishift.qb.") ||
326 Name.starts_with(
"pmulu.dq.") ||
327 Name.starts_with(
"por.") ||
328 Name.starts_with(
"prol.") ||
329 Name.starts_with(
"prolv.") ||
330 Name.starts_with(
"pror.") ||
331 Name.starts_with(
"prorv.") ||
332 Name.starts_with(
"pshuf.b.") ||
333 Name.starts_with(
"pshuf.d.") ||
334 Name.starts_with(
"pshufh.w.") ||
335 Name.starts_with(
"pshufl.w.") ||
336 Name.starts_with(
"psll.d") ||
337 Name.starts_with(
"psll.q") ||
338 Name.starts_with(
"psll.w") ||
339 Name.starts_with(
"pslli") ||
340 Name.starts_with(
"psllv") ||
341 Name.starts_with(
"psra.d") ||
342 Name.starts_with(
"psra.q") ||
343 Name.starts_with(
"psra.w") ||
344 Name.starts_with(
"psrai") ||
345 Name.starts_with(
"psrav") ||
346 Name.starts_with(
"psrl.d") ||
347 Name.starts_with(
"psrl.q") ||
348 Name.starts_with(
"psrl.w") ||
349 Name.starts_with(
"psrli") ||
350 Name.starts_with(
"psrlv") ||
351 Name.starts_with(
"psub.") ||
352 Name.starts_with(
"psubs.") ||
353 Name.starts_with(
"psubus.") ||
354 Name.starts_with(
"pternlog.") ||
355 Name.starts_with(
"punpckh") ||
356 Name.starts_with(
"punpckl") ||
357 Name.starts_with(
"pxor.") ||
358 Name.starts_with(
"shuf.f") ||
359 Name.starts_with(
"shuf.i") ||
360 Name.starts_with(
"shuf.p") ||
361 Name.starts_with(
"sqrt.p") ||
362 Name.starts_with(
"store.b.") ||
363 Name.starts_with(
"store.d.") ||
364 Name.starts_with(
"store.p") ||
365 Name.starts_with(
"store.q.") ||
366 Name.starts_with(
"store.w.") ||
367 Name ==
"store.ss" ||
368 Name.starts_with(
"storeu.") ||
369 Name.starts_with(
"sub.p") ||
370 Name.starts_with(
"ucmp.") ||
371 Name.starts_with(
"unpckh.") ||
372 Name.starts_with(
"unpckl.") ||
373 Name.starts_with(
"valign.") ||
374 Name ==
"vcvtph2ps.128" ||
375 Name ==
"vcvtph2ps.256" ||
376 Name.starts_with(
"vextract") ||
377 Name.starts_with(
"vfmadd.") ||
378 Name.starts_with(
"vfmaddsub.") ||
379 Name.starts_with(
"vfnmadd.") ||
380 Name.starts_with(
"vfnmsub.") ||
381 Name.starts_with(
"vpdpbusd.") ||
382 Name.starts_with(
"vpdpbusds.") ||
383 Name.starts_with(
"vpdpwssd.") ||
384 Name.starts_with(
"vpdpwssds.") ||
385 Name.starts_with(
"vpermi2var.") ||
386 Name.starts_with(
"vpermil.p") ||
387 Name.starts_with(
"vpermilvar.") ||
388 Name.starts_with(
"vpermt2var.") ||
389 Name.starts_with(
"vpmadd52") ||
390 Name.starts_with(
"vpshld.") ||
391 Name.starts_with(
"vpshldv.") ||
392 Name.starts_with(
"vpshrd.") ||
393 Name.starts_with(
"vpshrdv.") ||
394 Name.starts_with(
"vpshufbitqmb.") ||
395 Name.starts_with(
"xor."));
397 if (Name.consume_front(
"mask3."))
399 return (Name.starts_with(
"vfmadd.") ||
400 Name.starts_with(
"vfmaddsub.") ||
401 Name.starts_with(
"vfmsub.") ||
402 Name.starts_with(
"vfmsubadd.") ||
403 Name.starts_with(
"vfnmsub."));
405 if (Name.consume_front(
"maskz."))
407 return (Name.starts_with(
"pternlog.") ||
408 Name.starts_with(
"vfmadd.") ||
409 Name.starts_with(
"vfmaddsub.") ||
410 Name.starts_with(
"vpdpbusd.") ||
411 Name.starts_with(
"vpdpbusds.") ||
412 Name.starts_with(
"vpdpwssd.") ||
413 Name.starts_with(
"vpdpwssds.") ||
414 Name.starts_with(
"vpermt2var.") ||
415 Name.starts_with(
"vpmadd52") ||
416 Name.starts_with(
"vpshldv.") ||
417 Name.starts_with(
"vpshrdv."));
420 return (Name ==
"movntdqa" ||
421 Name ==
"pmul.dq.512" ||
422 Name ==
"pmulu.dq.512" ||
423 Name.starts_with(
"broadcastm") ||
424 Name.starts_with(
"cmp.p") ||
425 Name.starts_with(
"cvtb2mask.") ||
426 Name.starts_with(
"cvtd2mask.") ||
427 Name.starts_with(
"cvtmask2") ||
428 Name.starts_with(
"cvtq2mask.") ||
429 Name ==
"cvtusi2sd" ||
430 Name.starts_with(
"cvtw2mask.") ||
435 Name ==
"kortestc.w" ||
436 Name ==
"kortestz.w" ||
437 Name.starts_with(
"kunpck") ||
440 Name.starts_with(
"padds.") ||
441 Name.starts_with(
"pbroadcast") ||
442 Name.starts_with(
"pmulh.w") ||
443 Name.starts_with(
"pmulhu.w") ||
444 Name.starts_with(
"prol") ||
445 Name.starts_with(
"pror") ||
446 Name.starts_with(
"psll.dq") ||
447 Name.starts_with(
"psrl.dq") ||
448 Name.starts_with(
"psubs.") ||
449 Name.starts_with(
"ptestm") ||
450 Name.starts_with(
"ptestnm") ||
451 Name.starts_with(
"storent.") ||
452 Name.starts_with(
"vbroadcast.s") ||
453 Name.starts_with(
"vpshld.") ||
454 Name.starts_with(
"vpshrd."));
457 if (Name.consume_front(
"fma."))
458 return (Name.starts_with(
"vfmadd.") ||
459 Name.starts_with(
"vfmsub.") ||
460 Name.starts_with(
"vfmsubadd.") ||
461 Name.starts_with(
"vfnmadd.") ||
462 Name.starts_with(
"vfnmsub."));
464 if (Name.consume_front(
"fma4."))
465 return Name.starts_with(
"vfmadd.s");
467 if (Name.consume_front(
"sse."))
468 return (Name ==
"add.ss" ||
469 Name ==
"cvtsi2ss" ||
470 Name ==
"cvtsi642ss" ||
473 Name.starts_with(
"sqrt.p") ||
475 Name.starts_with(
"storeu.") ||
478 if (Name.consume_front(
"sse2."))
479 return (Name ==
"add.sd" ||
480 Name ==
"cvtdq2pd" ||
481 Name ==
"cvtdq2ps" ||
482 Name ==
"cvtps2pd" ||
483 Name ==
"cvtsi2sd" ||
484 Name ==
"cvtsi642sd" ||
485 Name ==
"cvtss2sd" ||
488 Name.starts_with(
"padds.") ||
489 Name.starts_with(
"paddus.") ||
490 Name.starts_with(
"pcmpeq.") ||
491 Name.starts_with(
"pcmpgt.") ||
497 Name ==
"pmulhu.w" ||
498 Name ==
"pmulu.dq" ||
499 Name.starts_with(
"pshuf") ||
500 Name.starts_with(
"psll.dq") ||
501 Name.starts_with(
"psrl.dq") ||
502 Name.starts_with(
"psubs.") ||
503 Name.starts_with(
"psubus.") ||
504 Name.starts_with(
"sqrt.p") ||
506 Name ==
"storel.dq" ||
507 Name.starts_with(
"storeu.") ||
510 if (Name.consume_front(
"sse41."))
511 return (Name.starts_with(
"blendp") ||
512 Name ==
"movntdqa" ||
522 Name.starts_with(
"pmovsx") ||
523 Name.starts_with(
"pmovzx") ||
526 if (Name.consume_front(
"sse42."))
527 return Name ==
"crc32.64.8";
529 if (Name.consume_front(
"sse4a."))
530 return Name.starts_with(
"movnt.");
532 if (Name.consume_front(
"ssse3."))
533 return (Name ==
"pabs.b.128" ||
534 Name ==
"pabs.d.128" ||
535 Name ==
"pabs.w.128");
537 if (Name.consume_front(
"xop."))
538 return (Name ==
"vpcmov" ||
539 Name ==
"vpcmov.256" ||
540 Name.starts_with(
"vpcom") ||
541 Name.starts_with(
"vprot"));
543 if (Name.consume_front(
"bmi."))
544 return (Name.starts_with(
"pdep.") ||
545 Name.starts_with(
"pext."));
547 return (Name ==
"addcarry.u32" ||
548 Name ==
"addcarry.u64" ||
549 Name ==
"addcarryx.u32" ||
550 Name ==
"addcarryx.u64" ||
551 Name ==
"subborrow.u32" ||
552 Name ==
"subborrow.u64" ||
553 Name.starts_with(
"vcvtph2ps."));
559 if (!Name.consume_front(
"x86."))
567 if (Name ==
"rdtscp") {
569 if (
F->getFunctionType()->getNumParams() == 0)
574 Intrinsic::x86_rdtscp);
581 if (Name.consume_front(
"sse41.ptest")) {
583 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
584 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
585 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
598 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
599 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
600 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
601 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
602 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
603 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
608 if (Name.consume_front(
"avx512.")) {
609 if (Name.consume_front(
"mask.cmp.")) {
612 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
613 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
614 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
615 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
616 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
617 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
621 }
else if (Name.starts_with(
"vpdpbusd.") ||
622 Name.starts_with(
"vpdpbusds.")) {
625 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
626 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
627 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
628 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
629 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
630 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
634 }
else if (Name.starts_with(
"vpdpwssd.") ||
635 Name.starts_with(
"vpdpwssds.")) {
638 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
639 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
640 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
641 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
642 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
643 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
651 if (Name.consume_front(
"avx2.")) {
652 if (Name.consume_front(
"vpdpb")) {
655 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
656 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
657 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
658 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
659 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
660 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
661 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
662 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
663 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
664 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
665 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
666 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
670 }
else if (Name.consume_front(
"vpdpw")) {
673 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
674 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
675 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
676 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
677 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
678 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
679 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
680 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
681 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
682 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
683 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
684 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
692 if (Name.consume_front(
"avx10.")) {
693 if (Name.consume_front(
"vpdpb")) {
696 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
697 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
698 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
699 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
700 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
701 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
705 }
else if (Name.consume_front(
"vpdpw")) {
707 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
708 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
709 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
710 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
711 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
712 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
720 if (Name.consume_front(
"avx512bf16.")) {
723 .
Case(
"cvtne2ps2bf16.128",
724 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
725 .
Case(
"cvtne2ps2bf16.256",
726 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
727 .
Case(
"cvtne2ps2bf16.512",
728 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
729 .
Case(
"mask.cvtneps2bf16.128",
730 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
731 .
Case(
"cvtneps2bf16.256",
732 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
733 .
Case(
"cvtneps2bf16.512",
734 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
741 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
742 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
743 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
750 if (Name.consume_front(
"xop.")) {
752 if (Name.starts_with(
"vpermil2")) {
755 auto Idx =
F->getFunctionType()->getParamType(2);
756 if (Idx->isFPOrFPVectorTy()) {
757 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
758 unsigned EltSize = Idx->getScalarSizeInBits();
759 if (EltSize == 64 && IdxSize == 128)
760 ID = Intrinsic::x86_xop_vpermil2pd;
761 else if (EltSize == 32 && IdxSize == 128)
762 ID = Intrinsic::x86_xop_vpermil2ps;
763 else if (EltSize == 64 && IdxSize == 256)
764 ID = Intrinsic::x86_xop_vpermil2pd_256;
766 ID = Intrinsic::x86_xop_vpermil2ps_256;
768 }
else if (
F->arg_size() == 2)
771 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
772 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
783 if (Name ==
"seh.recoverfp") {
785 Intrinsic::eh_recoverfp);
797 if (Name.starts_with(
"rbit")) {
800 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
804 if (Name ==
"thread.pointer") {
807 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
811 bool Neon = Name.consume_front(
"neon.");
816 if (Name.consume_front(
"bfdot.")) {
820 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
825 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
826 assert((OperandWidth == 64 || OperandWidth == 128) &&
827 "Unexpected operand width");
829 std::array<Type *, 2> Tys{
840 if (Name.consume_front(
"bfm")) {
842 if (Name.consume_back(
".v4f32.v16i8")) {
888 F->arg_begin()->getType());
892 if (Name.consume_front(
"vst")) {
894 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
898 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
899 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
902 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
903 Intrinsic::arm_neon_vst4lane};
905 auto fArgs =
F->getFunctionType()->params();
906 Type *Tys[] = {fArgs[0], fArgs[1]};
909 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
912 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
921 if (Name.consume_front(
"mve.")) {
923 if (Name ==
"vctp64") {
933 if (Name.starts_with(
"vrintn.v")) {
935 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
940 if (Name.consume_back(
".v4i1")) {
942 if (Name.consume_back(
".predicated.v2i64.v4i32"))
944 return Name ==
"mull.int" || Name ==
"vqdmull";
946 if (Name.consume_back(
".v2i64")) {
948 bool IsGather = Name.consume_front(
"vldr.gather.");
949 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
950 if (Name.consume_front(
"base.")) {
952 Name.consume_front(
"wb.");
955 return Name ==
"predicated.v2i64";
958 if (Name.consume_front(
"offset.predicated."))
959 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
960 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
973 if (Name.consume_front(
"cde.vcx")) {
975 if (Name.consume_back(
".predicated.v2i64.v4i1"))
977 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
978 Name ==
"3q" || Name ==
"3qa";
992 F->arg_begin()->getType());
998 .
Case(
"smax", Intrinsic::smax)
999 .
Case(
"smin", Intrinsic::smin)
1000 .
Case(
"umax", Intrinsic::umax)
1001 .
Case(
"umin", Intrinsic::umin)
1004 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1007 F->getReturnType());
1011 if (Name.starts_with(
"addp")) {
1013 if (
F->arg_size() != 2)
1016 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1018 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1024 if (Name.starts_with(
"bfcvt")) {
1030 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1037 if (Name.consume_front(
"sve.")) {
1039 if (Name.consume_front(
"bf")) {
1040 if (Name ==
"mmla") {
1041 Type *Tys[] = {
F->getReturnType(),
1042 std::next(
F->arg_begin())->getType()};
1044 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1047 if (Name.consume_back(
".lane")) {
1051 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1052 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1053 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1065 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1070 if (Name.consume_front(
"convert.from.svbool")) {
1073 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1076 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1081 if (Name.consume_front(
"convert.to.svbool")) {
1084 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1087 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1092 if (Name.consume_front(
"addqv")) {
1094 if (!
F->getReturnType()->isFPOrFPVectorTy())
1097 auto Args =
F->getFunctionType()->params();
1098 Type *Tys[] = {
F->getReturnType(), Args[1]};
1100 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1104 if (Name.consume_front(
"ld")) {
1106 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1107 if (LdRegex.
match(Name)) {
1113 "Expected 2 arguments for ld* intrinsic.");
1114 Type *PtrTy =
F->getArg(1)->getType();
1117 Intrinsic::aarch64_sve_ld2_sret,
1118 Intrinsic::aarch64_sve_ld3_sret,
1119 Intrinsic::aarch64_sve_ld4_sret,
1122 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1128 if (Name.consume_front(
"tuple.")) {
1130 if (Name.starts_with(
"get")) {
1132 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1134 F->getParent(), Intrinsic::vector_extract, Tys);
1138 if (Name.starts_with(
"set")) {
1140 auto Args =
F->getFunctionType()->params();
1141 Type *Tys[] = {Args[0], Args[2], Args[1]};
1143 F->getParent(), Intrinsic::vector_insert, Tys);
1147 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1148 if (CreateTupleRegex.
match(Name)) {
1150 auto Args =
F->getFunctionType()->params();
1151 Type *Tys[] = {
F->getReturnType(), Args[1]};
1153 F->getParent(), Intrinsic::vector_insert, Tys);
1159 if (Name.starts_with(
"rev.nxv")) {
1162 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1168 if (Name.consume_front(
"sme.")) {
1170 if (Name.consume_front(
"ftmopa.")) {
1175 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1176 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1194#define NVVM_TMA_G2S_MODES(M) \
1195 M(tile_1d, "tile.1d") \
1196 M(tile_2d, "tile.2d") \
1197 M(tile_3d, "tile.3d") \
1198 M(tile_4d, "tile.4d") \
1199 M(tile_5d, "tile.5d") \
1200 M(tile_gather4_2d, "tile.gather4.2d") \
1201 M(im2col_3d, "im2col.3d") \
1202 M(im2col_4d, "im2col.4d") \
1203 M(im2col_5d, "im2col.5d") \
1204 M(im2col_w_3d, "im2col.w.3d") \
1205 M(im2col_w_4d, "im2col.w.4d") \
1206 M(im2col_w_5d, "im2col.w.5d") \
1207 M(im2col_w_128_3d, "im2col.w.128.3d") \
1208 M(im2col_w_128_4d, "im2col.w.128.4d") \
1209 M(im2col_w_128_5d, "im2col.w.128.5d")
1221 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1224#define G2S_ID(ID_SUFFIX, NAME) \
1225 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1235 size_t NumParams =
F->getFunctionType()->getNumParams();
1239 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1246 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1247 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1248 "expected the i64 cache-hint after the multicast mask");
1249 Type *MaskTy = Params[MaskIdx];
1264 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1267#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1268 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1280 if (!
F->getFunctionType()
1281 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1297 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1302 size_t NumParams =
F->getFunctionType()->getNumParams();
1303 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1307 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1312 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1325 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1330 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1333 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1353 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1356 auto [RedOpName, ShapeName] = Name.split(
'.');
1361 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1362 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1363 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1364 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1365 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1366 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1367 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1368 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1374 if (Name.consume_front(
"mapa.shared.cluster"))
1375 if (
F->getReturnType()->getPointerAddressSpace() ==
1377 return Intrinsic::nvvm_mapa_shared_cluster;
1379 if (Name.consume_front(
"cp.async.bulk.")) {
1382 .
Case(
"shared.cta.to.cluster",
1383 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1387 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1397 if (!Name.consume_front(
"tcgen05.commit."))
1400 if (Name.consume_front(
"shared."))
1402 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1403 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1406 if (Name.consume_front(
"mc.shared.")) {
1408 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1412 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1413 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1422 if (
F->arg_size() != 2)
1425 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1426 Name.consume_front(
"tcgen05.alloc."))
1428 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1429 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1432 if (Name.consume_front(
"tcgen05.dealloc."))
1434 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1435 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1442 if (Name.consume_front(
"fma.rn."))
1444 .
Case(
"bf16", Intrinsic::nvvm_fma_rn_bf16)
1445 .
Case(
"bf16x2", Intrinsic::nvvm_fma_rn_bf16x2)
1446 .
Case(
"relu.bf16", Intrinsic::nvvm_fma_rn_relu_bf16)
1447 .
Case(
"relu.bf16x2", Intrinsic::nvvm_fma_rn_relu_bf16x2)
1450 if (Name.consume_front(
"fmax."))
1452 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1453 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1454 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1455 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1456 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1457 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1458 .
Case(
"ftz.nan.xorsign.abs.bf16",
1459 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1460 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1461 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1462 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1463 .
Case(
"ftz.xorsign.abs.bf16x2",
1464 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1465 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1466 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1467 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1468 .
Case(
"nan.xorsign.abs.bf16x2",
1469 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1470 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1471 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1474 if (Name.consume_front(
"fmin."))
1476 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1477 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1478 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1479 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1480 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1481 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1482 .
Case(
"ftz.nan.xorsign.abs.bf16",
1483 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1484 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1485 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1486 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1487 .
Case(
"ftz.xorsign.abs.bf16x2",
1488 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1489 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1490 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1491 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1492 .
Case(
"nan.xorsign.abs.bf16x2",
1493 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1494 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1495 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1498 if (Name.consume_front(
"neg."))
1500 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1501 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1510 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1515 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1518 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1521 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1522 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1530 {Intrinsic::nvvm_fadd, Intrinsic::nvvm_fadd_sat},
1531 {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat}};
1533 {Intrinsic::nvvm_fmul, Intrinsic::nvvm_fmul_sat},
1534 {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat}};
1536static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1538 auto [Modifiers,
Type] = Name.rsplit(
'.');
1540 return std::nullopt;
1550 return std::nullopt;
1552 StringRef Rest = Modifiers.drop_front(2);
1556 return std::nullopt;
1558 return std::make_pair(IIDs[IsFTZ][IsSat], *
RoundingMode);
1562 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1565 return Intrinsic::nvvm_mbarrier_init;
1569 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1570 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1571 Name.consume_front(
"param");
1575 if (!Name.consume_front(
"vp."))
1604 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1605 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1612 if (!Name.consume_front(
"vp."))
1632 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1633 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1638 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1650 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1661 if (Name.starts_with(
"to.fp16")) {
1665 FuncTy->getReturnType());
1668 if (Name.starts_with(
"from.fp16")) {
1672 FuncTy->getReturnType());
1682 if (Defaults.empty())
1685 unsigned FullArgCount = FirstDefault + Defaults.size();
1688 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1691 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1693 NumMissingTrailingParams))
1696 return FullArgCount;
1703 unsigned FullArgCount =
1705 if (FullArgCount == 0)
1711 "total number of default args does not match intrinsic signature");
1716 bool CanUpgradeDebugIntrinsicsToRecords) {
1717 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1722 if (!Name.consume_front(
"llvm.") || Name.empty())
1728 bool IsArm = Name.consume_front(
"arm.");
1729 if (IsArm || Name.consume_front(
"aarch64.")) {
1735 if (Name.consume_front(
"amdgcn.")) {
1736 if (Name ==
"alignbit") {
1739 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1743 if (Name.consume_front(
"atomic.")) {
1744 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1745 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1754 if (Name.starts_with(
"addrspacecast.nonnull")) {
1761 switch (
F->getIntrinsicID()) {
1765 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1766 if (
F->arg_size() == 7) {
1771 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1772 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1773 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1774 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1775 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1776 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1777 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1778 if (
F->arg_size() == 8) {
1785 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1786 Name.consume_front(
"flat.atomic.")) {
1787 if (Name.starts_with(
"fadd") ||
1789 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1790 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1798 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1803 if (Name.starts_with(
"ldexp.")) {
1806 F->getParent(), Intrinsic::ldexp,
1807 {F->getReturnType(), F->getArg(1)->getType()});
1816 if (
F->arg_size() == 1) {
1817 if (Name.consume_front(
"convert.")) {
1831 F->arg_begin()->getType());
1837 if (Name ==
"coro.end" &&
1838 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1839 CoroEndID = Intrinsic::coro_end;
1840 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1841 CoroEndID = Intrinsic::coro_end_async;
1852 if (Name.consume_front(
"dbg.")) {
1854 if (CanUpgradeDebugIntrinsicsToRecords) {
1855 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1856 Name ==
"declare" || Name ==
"label") {
1865 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1868 Intrinsic::dbg_value);
1875 if (Name.consume_front(
"experimental.vector.")) {
1881 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1882 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1883 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1884 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1885 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1887 Intrinsic::vector_partial_reduce_add)
1890 const auto *FT =
F->getFunctionType();
1892 if (ID == Intrinsic::vector_extract ||
1893 ID == Intrinsic::vector_interleave2)
1896 if (ID != Intrinsic::vector_interleave2)
1898 if (ID == Intrinsic::vector_insert ||
1899 ID == Intrinsic::vector_partial_reduce_add)
1907 if (Name.consume_front(
"reduce.")) {
1909 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1910 if (R.match(Name, &
Groups))
1912 .
Case(
"add", Intrinsic::vector_reduce_add)
1913 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1914 .
Case(
"and", Intrinsic::vector_reduce_and)
1915 .
Case(
"or", Intrinsic::vector_reduce_or)
1916 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1917 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1918 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1919 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1920 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1921 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1922 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1927 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1932 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1933 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1938 auto Args =
F->getFunctionType()->params();
1940 {Args[V2 ? 1 : 0]});
1946 if (Name.consume_front(
"splice"))
1950 if (Name.consume_front(
"experimental.stepvector.")) {
1954 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1959 if (Name.starts_with(
"flt.rounds")) {
1962 Intrinsic::get_rounding);
1967 if (Name.starts_with(
"invariant.group.barrier")) {
1969 auto Args =
F->getFunctionType()->params();
1970 Type* ObjectPtr[1] = {Args[0]};
1973 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
1978 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
1979 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
1980 if (IsLifetimeStart || IsLifetimeEnd) {
1981 if (
F->arg_size() == 2) {
1982 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
1983 : Intrinsic::lifetime_end;
1988 F->getArg(1)->getType());
1990 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2010 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2011 .StartsWith(
"memmove.", Intrinsic::memmove)
2013 if (
F->arg_size() == 5) {
2017 F->getFunctionType()->params().slice(0, 3);
2023 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2026 const auto *FT =
F->getFunctionType();
2027 Type *ParamTypes[2] = {
2028 FT->getParamType(0),
2032 Intrinsic::memset, ParamTypes);
2038 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2039 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2040 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2041 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2043 if (MaskedID &&
F->arg_size() == 4) {
2045 if (MaskedID == Intrinsic::masked_load ||
2046 MaskedID == Intrinsic::masked_gather) {
2048 F->getParent(), MaskedID,
2049 {F->getReturnType(), F->getArg(0)->getType()});
2053 F->getParent(), MaskedID,
2054 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2060 if (Name.consume_front(
"nvvm.")) {
2062 if (
F->arg_size() == 1) {
2065 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2066 .Case(
"clz.i", Intrinsic::ctlz)
2067 .
Case(
"popc.i", Intrinsic::ctpop)
2071 {F->getReturnType()});
2074 }
else if (
F->arg_size() == 2) {
2077 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2078 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2079 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2080 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2081 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2082 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2086 {F->getReturnType()});
2123 F->getParent(), IID,
F->getReturnType(),
2124 F->getFunctionType()->params());
2135 {F->getArg(0)->getType()});
2183 F->getArg(0)->getType());
2191 bool Expand =
false;
2192 if (Name.consume_front(
"abs."))
2195 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2196 else if (Name.consume_front(
"fabs."))
2198 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2199 else if (Name.consume_front(
"add."))
2202 else if (Name.consume_front(
"mul."))
2205 else if (Name.consume_front(
"ex2.approx."))
2208 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2209 else if (Name.consume_front(
"atomic.load."))
2218 else if (Name.consume_front(
"atomic."))
2233 else if (Name.consume_front(
"bitcast."))
2236 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2237 else if (Name.consume_front(
"rotate."))
2239 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2240 else if (Name.consume_front(
"ptr.gen.to."))
2243 else if (Name.consume_front(
"ptr."))
2246 else if (Name.consume_front(
"ldg.global."))
2248 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2249 Name.starts_with(
"p."));
2252 .
Case(
"barrier0",
true)
2253 .
Case(
"barrier.n",
true)
2254 .
Case(
"barrier.sync.cnt",
true)
2255 .
Case(
"barrier.sync",
true)
2256 .
Case(
"barrier",
true)
2257 .
Case(
"bar.sync",
true)
2258 .
Case(
"barrier0.popc",
true)
2259 .
Case(
"barrier0.and",
true)
2260 .
Case(
"barrier0.or",
true)
2261 .
Case(
"clz.ll",
true)
2262 .
Case(
"popc.ll",
true)
2264 .
Case(
"swap.lo.hi.b64",
true)
2265 .
Case(
"tanh.approx.f32",
true)
2277 if (Name.starts_with(
"objectsize.")) {
2278 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2279 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2282 Intrinsic::objectsize, Tys);
2289 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2292 F->getParent(), Intrinsic::ptr_annotation,
2293 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2299 if (Name.consume_front(
"riscv.")) {
2302 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2303 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2304 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2305 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2308 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2321 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2322 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2331 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2332 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2333 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2334 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2339 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2348 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2350 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2359 if (Name ==
"stackprotectorcheck") {
2363 if (Name.starts_with(
"strip.invariant.group")) {
2368 F->getParent(), Intrinsic::launder_invariant_group,
2369 F->getReturnType());
2375 if (Name ==
"thread.pointer") {
2377 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2383 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2386 F->getParent(), Intrinsic::var_annotation,
2387 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2390 if (Name.consume_front(
"vector.splice")) {
2391 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2401 if (Name.consume_front(
"wasm.")) {
2404 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2405 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2406 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2411 F->getReturnType());
2415 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2417 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2419 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2438 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2448 std::string
Name =
F->getName().str();
2451 Name,
F->getParent());
2462 if (Result != std::nullopt) {
2479 bool CanUpgradeDebugIntrinsicsToRecords) {
2499 GV->
getName() ==
"llvm.global_dtors")) ||
2513 unsigned N =
Init->getNumOperands();
2514 std::vector<Constant *> NewCtors(
N);
2515 for (
unsigned i = 0; i !=
N; ++i) {
2518 Ctor->getAggregateElement(1),
2532 unsigned NumElts = ResultTy->getNumElements() * 8;
2536 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2546 for (
unsigned l = 0; l != NumElts; l += 16)
2547 for (
unsigned i = 0; i != 16; ++i) {
2548 unsigned Idx = NumElts + i - Shift;
2550 Idx -= NumElts - 16;
2551 Idxs[l + i] = Idx + l;
2554 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2558 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2566 unsigned NumElts = ResultTy->getNumElements() * 8;
2570 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2580 for (
unsigned l = 0; l != NumElts; l += 16)
2581 for (
unsigned i = 0; i != 16; ++i) {
2582 unsigned Idx = i + Shift;
2584 Idx += NumElts - 16;
2585 Idxs[l + i] = Idx + l;
2588 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2592 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2600 Mask = Builder.CreateBitCast(Mask, MaskTy);
2606 for (
unsigned i = 0; i != NumElts; ++i)
2608 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2619 if (
C->isAllOnesValue())
2624 return Builder.CreateSelect(Mask, Op0, Op1);
2631 if (
C->isAllOnesValue())
2635 Mask->getType()->getIntegerBitWidth());
2636 Mask = Builder.CreateBitCast(Mask, MaskTy);
2637 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2638 return Builder.CreateSelect(Mask, Op0, Op1);
2651 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2652 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2657 ShiftVal &= (NumElts - 1);
2666 if (ShiftVal > 16) {
2674 for (
unsigned l = 0; l < NumElts; l += 16) {
2675 for (
unsigned i = 0; i != 16; ++i) {
2676 unsigned Idx = ShiftVal + i;
2677 if (!IsVALIGN && Idx >= 16)
2678 Idx += NumElts - 16;
2679 Indices[l + i] = Idx + l;
2684 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2690 bool ZeroMask,
bool IndexForm) {
2693 unsigned EltWidth = Ty->getScalarSizeInBits();
2694 bool IsFloat = Ty->isFPOrFPVectorTy();
2696 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2697 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2698 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2699 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2700 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2701 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2702 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2703 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2704 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2705 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2706 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2707 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2708 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2709 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2710 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2711 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2712 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2713 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2714 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2715 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2716 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2717 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2718 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2719 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2720 else if (VecWidth == 128 && EltWidth == 16)
2721 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2722 else if (VecWidth == 256 && EltWidth == 16)
2723 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2724 else if (VecWidth == 512 && EltWidth == 16)
2725 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2726 else if (VecWidth == 128 && EltWidth == 8)
2727 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2728 else if (VecWidth == 256 && EltWidth == 8)
2729 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2730 else if (VecWidth == 512 && EltWidth == 8)
2731 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2742 Value *V = Builder.CreateIntrinsic(IID, Args);
2754 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2765 bool IsRotateRight) {
2775 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2776 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2779 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2780 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2825 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2830 bool IsShiftRight,
bool ZeroMask) {
2844 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2845 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2848 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2849 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2864 const Align Alignment =
2866 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2871 if (
C->isAllOnesValue())
2872 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2877 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2883 const Align Alignment =
2892 if (
C->isAllOnesValue())
2893 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2898 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2904 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2905 {Op0, Builder.getInt1(
false)});
2920 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2921 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2922 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2923 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2924 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2927 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2928 LHS = Builder.CreateAnd(
LHS, Mask);
2929 RHS = Builder.CreateAnd(
RHS, Mask);
2946 if (!
C || !
C->isAllOnesValue())
2947 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2952 for (
unsigned i = 0; i != NumElts; ++i)
2954 for (
unsigned i = NumElts; i != 8; ++i)
2955 Indices[i] = NumElts + i % NumElts;
2956 Vec = Builder.CreateShuffleVector(Vec,
2960 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2964 unsigned CC,
bool Signed) {
2972 }
else if (CC == 7) {
3008 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3009 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3011 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3012 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3021 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3027 Name = Name.substr(12);
3032 if (Name.starts_with(
"max.p")) {
3033 if (VecWidth == 128 && EltWidth == 32)
3034 IID = Intrinsic::x86_sse_max_ps;
3035 else if (VecWidth == 128 && EltWidth == 64)
3036 IID = Intrinsic::x86_sse2_max_pd;
3037 else if (VecWidth == 256 && EltWidth == 32)
3038 IID = Intrinsic::x86_avx_max_ps_256;
3039 else if (VecWidth == 256 && EltWidth == 64)
3040 IID = Intrinsic::x86_avx_max_pd_256;
3043 }
else if (Name.starts_with(
"min.p")) {
3044 if (VecWidth == 128 && EltWidth == 32)
3045 IID = Intrinsic::x86_sse_min_ps;
3046 else if (VecWidth == 128 && EltWidth == 64)
3047 IID = Intrinsic::x86_sse2_min_pd;
3048 else if (VecWidth == 256 && EltWidth == 32)
3049 IID = Intrinsic::x86_avx_min_ps_256;
3050 else if (VecWidth == 256 && EltWidth == 64)
3051 IID = Intrinsic::x86_avx_min_pd_256;
3054 }
else if (Name.starts_with(
"pshuf.b.")) {
3055 if (VecWidth == 128)
3056 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3057 else if (VecWidth == 256)
3058 IID = Intrinsic::x86_avx2_pshuf_b;
3059 else if (VecWidth == 512)
3060 IID = Intrinsic::x86_avx512_pshuf_b_512;
3063 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3064 if (VecWidth == 128)
3065 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3066 else if (VecWidth == 256)
3067 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3068 else if (VecWidth == 512)
3069 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3072 }
else if (Name.starts_with(
"pmulh.w")) {
3073 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3074 "Unexpected intrinsic");
3077 }
else if (Name.starts_with(
"pmulhu.w")) {
3078 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3079 "Unexpected intrinsic");
3082 }
else if (Name.starts_with(
"pmaddw.d.")) {
3083 if (VecWidth == 128)
3084 IID = Intrinsic::x86_sse2_pmadd_wd;
3085 else if (VecWidth == 256)
3086 IID = Intrinsic::x86_avx2_pmadd_wd;
3087 else if (VecWidth == 512)
3088 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3091 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3092 if (VecWidth == 128)
3093 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3094 else if (VecWidth == 256)
3095 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3096 else if (VecWidth == 512)
3097 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3100 }
else if (Name.starts_with(
"packsswb.")) {
3101 if (VecWidth == 128)
3102 IID = Intrinsic::x86_sse2_packsswb_128;
3103 else if (VecWidth == 256)
3104 IID = Intrinsic::x86_avx2_packsswb;
3105 else if (VecWidth == 512)
3106 IID = Intrinsic::x86_avx512_packsswb_512;
3109 }
else if (Name.starts_with(
"packssdw.")) {
3110 if (VecWidth == 128)
3111 IID = Intrinsic::x86_sse2_packssdw_128;
3112 else if (VecWidth == 256)
3113 IID = Intrinsic::x86_avx2_packssdw;
3114 else if (VecWidth == 512)
3115 IID = Intrinsic::x86_avx512_packssdw_512;
3118 }
else if (Name.starts_with(
"packuswb.")) {
3119 if (VecWidth == 128)
3120 IID = Intrinsic::x86_sse2_packuswb_128;
3121 else if (VecWidth == 256)
3122 IID = Intrinsic::x86_avx2_packuswb;
3123 else if (VecWidth == 512)
3124 IID = Intrinsic::x86_avx512_packuswb_512;
3127 }
else if (Name.starts_with(
"packusdw.")) {
3128 if (VecWidth == 128)
3129 IID = Intrinsic::x86_sse41_packusdw;
3130 else if (VecWidth == 256)
3131 IID = Intrinsic::x86_avx2_packusdw;
3132 else if (VecWidth == 512)
3133 IID = Intrinsic::x86_avx512_packusdw_512;
3136 }
else if (Name.starts_with(
"vpermilvar.")) {
3137 if (VecWidth == 128 && EltWidth == 32)
3138 IID = Intrinsic::x86_avx_vpermilvar_ps;
3139 else if (VecWidth == 128 && EltWidth == 64)
3140 IID = Intrinsic::x86_avx_vpermilvar_pd;
3141 else if (VecWidth == 256 && EltWidth == 32)
3142 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3143 else if (VecWidth == 256 && EltWidth == 64)
3144 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3145 else if (VecWidth == 512 && EltWidth == 32)
3146 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3147 else if (VecWidth == 512 && EltWidth == 64)
3148 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3151 }
else if (Name ==
"cvtpd2dq.256") {
3152 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3153 }
else if (Name ==
"cvtpd2ps.256") {
3154 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3155 }
else if (Name ==
"cvttpd2dq.256") {
3156 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3157 }
else if (Name ==
"cvttps2dq.128") {
3158 IID = Intrinsic::x86_sse2_cvttps2dq;
3159 }
else if (Name ==
"cvttps2dq.256") {
3160 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3161 }
else if (Name.starts_with(
"permvar.")) {
3163 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3164 IID = Intrinsic::x86_avx2_permps;
3165 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3166 IID = Intrinsic::x86_avx2_permd;
3167 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3168 IID = Intrinsic::x86_avx512_permvar_df_256;
3169 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3170 IID = Intrinsic::x86_avx512_permvar_di_256;
3171 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3172 IID = Intrinsic::x86_avx512_permvar_sf_512;
3173 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3174 IID = Intrinsic::x86_avx512_permvar_si_512;
3175 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3176 IID = Intrinsic::x86_avx512_permvar_df_512;
3177 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3178 IID = Intrinsic::x86_avx512_permvar_di_512;
3179 else if (VecWidth == 128 && EltWidth == 16)
3180 IID = Intrinsic::x86_avx512_permvar_hi_128;
3181 else if (VecWidth == 256 && EltWidth == 16)
3182 IID = Intrinsic::x86_avx512_permvar_hi_256;
3183 else if (VecWidth == 512 && EltWidth == 16)
3184 IID = Intrinsic::x86_avx512_permvar_hi_512;
3185 else if (VecWidth == 128 && EltWidth == 8)
3186 IID = Intrinsic::x86_avx512_permvar_qi_128;
3187 else if (VecWidth == 256 && EltWidth == 8)
3188 IID = Intrinsic::x86_avx512_permvar_qi_256;
3189 else if (VecWidth == 512 && EltWidth == 8)
3190 IID = Intrinsic::x86_avx512_permvar_qi_512;
3193 }
else if (Name.starts_with(
"dbpsadbw.")) {
3194 if (VecWidth == 128)
3195 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3196 else if (VecWidth == 256)
3197 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3198 else if (VecWidth == 512)
3199 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3202 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3203 if (VecWidth == 128)
3204 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3205 else if (VecWidth == 256)
3206 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3207 else if (VecWidth == 512)
3208 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3211 }
else if (Name.starts_with(
"conflict.")) {
3212 if (Name[9] ==
'd' && VecWidth == 128)
3213 IID = Intrinsic::x86_avx512_conflict_d_128;
3214 else if (Name[9] ==
'd' && VecWidth == 256)
3215 IID = Intrinsic::x86_avx512_conflict_d_256;
3216 else if (Name[9] ==
'd' && VecWidth == 512)
3217 IID = Intrinsic::x86_avx512_conflict_d_512;
3218 else if (Name[9] ==
'q' && VecWidth == 128)
3219 IID = Intrinsic::x86_avx512_conflict_q_128;
3220 else if (Name[9] ==
'q' && VecWidth == 256)
3221 IID = Intrinsic::x86_avx512_conflict_q_256;
3222 else if (Name[9] ==
'q' && VecWidth == 512)
3223 IID = Intrinsic::x86_avx512_conflict_q_512;
3226 }
else if (Name.starts_with(
"pavg.")) {
3227 if (Name[5] ==
'b' && VecWidth == 128)
3228 IID = Intrinsic::x86_sse2_pavg_b;
3229 else if (Name[5] ==
'b' && VecWidth == 256)
3230 IID = Intrinsic::x86_avx2_pavg_b;
3231 else if (Name[5] ==
'b' && VecWidth == 512)
3232 IID = Intrinsic::x86_avx512_pavg_b_512;
3233 else if (Name[5] ==
'w' && VecWidth == 128)
3234 IID = Intrinsic::x86_sse2_pavg_w;
3235 else if (Name[5] ==
'w' && VecWidth == 256)
3236 IID = Intrinsic::x86_avx2_pavg_w;
3237 else if (Name[5] ==
'w' && VecWidth == 512)
3238 IID = Intrinsic::x86_avx512_pavg_w_512;
3247 Rep = Builder.CreateIntrinsic(IID, Args);
3258 if (AsmStr->find(
"mov\tfp") == 0 &&
3259 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3260 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3261 AsmStr->replace(Pos, 1,
";");
3269 assert(Result &&
"unsupported nvvm.add.*/nvvm.mul.* intrinsic");
3272 return Builder.CreateIntrinsic(
3274 {A, CI->getArgOperand(1),
3275 Builder.getInt32(static_cast<int>(RoundingMode))});
3280 Value *Rep =
nullptr;
3282 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3284 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3285 {Arg, Builder.getTrue()},
3287 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3288 Type *Ty = (Name ==
"abs.bf16")
3292 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3293 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3294 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3295 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3296 : Intrinsic::nvvm_fabs;
3297 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3298 }
else if (Name.consume_front(
"add.")) {
3301 }
else if (Name.consume_front(
"mul.")) {
3304 }
else if (Name.consume_front(
"ex2.approx.")) {
3306 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3307 : Intrinsic::nvvm_ex2_approx;
3308 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3309 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3310 Name.starts_with(
"atomic.load.add.f64.p")) {
3313 Rep = Builder.CreateAtomicRMW(
3317 LLVMContext::MD_atomic_ignore_denormal_mode,
3322 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3323 Name.starts_with(
"atomic.load.dec.32.p")) {
3328 Rep = Builder.CreateAtomicRMW(
3332 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3338 Op.contains(
".cta.") ?
"block" :
"");
3339 if (
Op.starts_with(
"cas.")) {
3341 Value *Pair = Builder.CreateAtomicCmpXchg(
3344 Rep = Builder.CreateExtractValue(Pair, 0);
3362 "unexpected nvvm scoped atomic intrinsic");
3363 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3367 LLVMContext::MD_atomic_ignore_denormal_mode,
3371 }
else if (Name ==
"clz.ll") {
3374 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3375 {Arg, Builder.getFalse()},
3377 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3378 }
else if (Name ==
"popc.ll") {
3382 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3383 Arg,
nullptr,
"ctpop");
3384 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3385 }
else if (Name ==
"h2f") {
3387 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3388 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3389 }
else if (Name.consume_front(
"bitcast.") &&
3390 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3393 }
else if (Name ==
"rotate.b32") {
3396 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3397 {Arg, Arg, ShiftAmt});
3398 }
else if (Name ==
"rotate.b64") {
3402 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3403 {Arg, Arg, ZExtShiftAmt});
3404 }
else if (Name ==
"rotate.right.b64") {
3408 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3409 {Arg, Arg, ZExtShiftAmt});
3410 }
else if (Name ==
"swap.lo.hi.b64") {
3413 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3414 {Arg, Arg, Builder.getInt64(32)});
3415 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3418 Name.starts_with(
".to.gen"))) {
3420 }
else if (Name.consume_front(
"ldg.global")) {
3424 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3427 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3429 }
else if (Name ==
"tanh.approx.f32") {
3433 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3435 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3437 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3438 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3440 }
else if (Name ==
"barrier") {
3441 Rep = Builder.CreateIntrinsic(
3442 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3444 }
else if (Name ==
"barrier.sync") {
3445 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3447 }
else if (Name ==
"barrier.sync.cnt") {
3448 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3450 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3451 Name ==
"barrier0.or") {
3453 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3457 .
Case(
"barrier0.popc",
3458 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3459 .
Case(
"barrier0.and",
3460 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3461 .
Case(
"barrier0.or",
3462 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3463 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3464 Rep = Builder.CreateZExt(Bar, CI->
getType());
3478 ? Builder.CreateBitCast(Arg, NewType)
3481 Rep = Builder.CreateCall(NewFn, Args);
3482 if (
F->getReturnType()->isIntegerTy())
3483 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3493 Value *Rep =
nullptr;
3495 if (Name.starts_with(
"sse4a.movnt.")) {
3507 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3510 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3511 }
else if (Name.starts_with(
"avx.movnt.") ||
3512 Name.starts_with(
"avx512.storent.")) {
3524 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3525 }
else if (Name ==
"sse2.storel.dq") {
3530 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3531 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3532 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3533 }
else if (Name.starts_with(
"sse.storeu.") ||
3534 Name.starts_with(
"sse2.storeu.") ||
3535 Name.starts_with(
"avx.storeu.")) {
3538 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3539 }
else if (Name ==
"avx512.mask.store.ss") {
3543 }
else if (Name.starts_with(
"avx512.mask.store")) {
3545 bool Aligned = Name[17] !=
'u';
3548 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3551 bool CmpEq = Name[9] ==
'e';
3554 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3555 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3562 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3563 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3565 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3566 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3567 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3568 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3569 Name.starts_with(
"sse2.sqrt.p") ||
3570 Name.starts_with(
"sse.sqrt.p")) {
3571 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3572 {CI->getArgOperand(0)});
3573 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3577 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3578 : Intrinsic::x86_avx512_sqrt_pd_512;
3581 Rep = Builder.CreateIntrinsic(IID, Args);
3583 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3584 {CI->getArgOperand(0)});
3588 }
else if (Name.starts_with(
"avx512.ptestm") ||
3589 Name.starts_with(
"avx512.ptestnm")) {
3593 Rep = Builder.CreateAnd(Op0, Op1);
3599 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3601 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3604 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3607 }
else if (Name.starts_with(
"avx512.kunpck")) {
3612 for (
unsigned i = 0; i != NumElts; ++i)
3621 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3622 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3623 }
else if (Name ==
"avx512.kand.w") {
3626 Rep = Builder.CreateAnd(
LHS,
RHS);
3627 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3628 }
else if (Name ==
"avx512.kandn.w") {
3631 LHS = Builder.CreateNot(
LHS);
3632 Rep = Builder.CreateAnd(
LHS,
RHS);
3633 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3634 }
else if (Name ==
"avx512.kor.w") {
3637 Rep = Builder.CreateOr(
LHS,
RHS);
3638 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3639 }
else if (Name ==
"avx512.kxor.w") {
3642 Rep = Builder.CreateXor(
LHS,
RHS);
3643 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3644 }
else if (Name ==
"avx512.kxnor.w") {
3647 LHS = Builder.CreateNot(
LHS);
3648 Rep = Builder.CreateXor(
LHS,
RHS);
3649 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3650 }
else if (Name ==
"avx512.knot.w") {
3652 Rep = Builder.CreateNot(Rep);
3653 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3654 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3657 Rep = Builder.CreateOr(
LHS,
RHS);
3658 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3660 if (Name[14] ==
'c')
3664 Rep = Builder.CreateICmpEQ(Rep,
C);
3665 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3666 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3667 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3668 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3669 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3672 ConstantInt::get(I32Ty, 0));
3674 ConstantInt::get(I32Ty, 0));
3676 if (Name.contains(
".add."))
3677 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3678 else if (Name.contains(
".sub."))
3679 EltOp = Builder.CreateFSub(Elt0, Elt1);
3680 else if (Name.contains(
".mul."))
3681 EltOp = Builder.CreateFMul(Elt0, Elt1);
3683 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3684 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3685 ConstantInt::get(I32Ty, 0));
3686 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3688 bool CmpEq = Name[16] ==
'e';
3690 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3692 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3699 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3702 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3705 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3712 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3714 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3715 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3717 if (VecWidth == 128 && EltWidth == 32)
3718 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3719 else if (VecWidth == 256 && EltWidth == 32)
3720 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3721 else if (VecWidth == 512 && EltWidth == 32)
3722 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3723 else if (VecWidth == 128 && EltWidth == 64)
3724 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3725 else if (VecWidth == 256 && EltWidth == 64)
3726 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3727 else if (VecWidth == 512 && EltWidth == 64)
3728 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3735 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3738 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3739 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3741 if (VecWidth == 128 && EltWidth == 32)
3742 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3743 else if (VecWidth == 256 && EltWidth == 32)
3744 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3745 else if (VecWidth == 512 && EltWidth == 32)
3746 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3747 else if (VecWidth == 128 && EltWidth == 64)
3748 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3749 else if (VecWidth == 256 && EltWidth == 64)
3750 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3751 else if (VecWidth == 512 && EltWidth == 64)
3752 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3757 if (VecWidth == 512)
3759 Args.push_back(Mask);
3761 Rep = Builder.CreateIntrinsic(IID, Args);
3762 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3766 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3769 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3770 Name.starts_with(
"avx512.cvtw2mask.") ||
3771 Name.starts_with(
"avx512.cvtd2mask.") ||
3772 Name.starts_with(
"avx512.cvtq2mask.")) {
3777 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3778 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3779 Name.starts_with(
"avx512.mask.pabs")) {
3781 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3782 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3783 Name.starts_with(
"avx512.mask.pmaxs")) {
3785 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3786 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3787 Name.starts_with(
"avx512.mask.pmaxu")) {
3789 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3790 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3791 Name.starts_with(
"avx512.mask.pmins")) {
3793 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3794 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3795 Name.starts_with(
"avx512.mask.pminu")) {
3797 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3798 Name.starts_with(
"avx512.pmulh.w")) {
3800 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3801 Name.starts_with(
"avx512.pmulhu.w")) {
3803 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3804 Name ==
"avx512.pmulu.dq.512" ||
3805 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3807 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3808 Name ==
"avx512.pmul.dq.512" ||
3809 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3811 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3812 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3817 }
else if (Name ==
"avx512.cvtusi2sd") {
3822 }
else if (Name ==
"sse2.cvtss2sd") {
3824 Rep = Builder.CreateFPExt(
3827 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3828 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3829 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3830 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3831 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3832 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3833 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3834 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3835 Name ==
"avx512.mask.cvtqq2ps.256" ||
3836 Name ==
"avx512.mask.cvtqq2ps.512" ||
3837 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3838 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3839 Name ==
"avx.cvt.ps2.pd.256" ||
3840 Name ==
"avx512.mask.cvtps2pd.128" ||
3841 Name ==
"avx512.mask.cvtps2pd.256") {
3846 unsigned NumDstElts = DstTy->getNumElements();
3847 if (NumDstElts < SrcTy->getNumElements()) {
3848 assert(NumDstElts == 2 &&
"Unexpected vector size");
3849 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3852 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3853 bool IsUnsigned = Name.contains(
"cvtu");
3855 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3859 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3860 : Intrinsic::x86_avx512_sitofp_round;
3861 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3864 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3865 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3871 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3872 Name.starts_with(
"vcvtph2ps.")) {
3876 unsigned NumDstElts = DstTy->getNumElements();
3877 if (NumDstElts != SrcTy->getNumElements()) {
3878 assert(NumDstElts == 4 &&
"Unexpected vector size");
3879 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3881 Rep = Builder.CreateBitCast(
3883 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3887 }
else if (Name.starts_with(
"avx512.mask.load")) {
3889 bool Aligned = Name[16] !=
'u';
3892 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3896 ResultTy->getNumElements());
3897 Rep = Builder.CreateIntrinsic(
3898 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3900 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3906 Rep = Builder.CreateIntrinsic(
3907 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3909 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3910 Name.starts_with(
"avx512.mask.expand.")) {
3914 ResultTy->getNumElements());
3916 bool IsCompress = Name[12] ==
'c';
3917 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3918 : Intrinsic::x86_avx512_mask_expand;
3919 Rep = Builder.CreateIntrinsic(
3921 }
else if (Name.starts_with(
"xop.vpcom")) {
3923 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3924 Name.ends_with(
"uq"))
3926 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3927 Name.ends_with(
"d") || Name.ends_with(
"q"))
3936 Name = Name.substr(9);
3937 if (Name.starts_with(
"lt"))
3939 else if (Name.starts_with(
"le"))
3941 else if (Name.starts_with(
"gt"))
3943 else if (Name.starts_with(
"ge"))
3945 else if (Name.starts_with(
"eq"))
3947 else if (Name.starts_with(
"ne"))
3949 else if (Name.starts_with(
"false"))
3951 else if (Name.starts_with(
"true"))
3958 }
else if (Name.starts_with(
"xop.vpcmov")) {
3960 Value *NotSel = Builder.CreateNot(Sel);
3963 Rep = Builder.CreateOr(Sel0, Sel1);
3964 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
3965 Name.starts_with(
"avx512.mask.prol")) {
3967 }
else if (Name.starts_with(
"avx512.pror") ||
3968 Name.starts_with(
"avx512.mask.pror")) {
3970 }
else if (Name.starts_with(
"avx512.vpshld.") ||
3971 Name.starts_with(
"avx512.mask.vpshld") ||
3972 Name.starts_with(
"avx512.maskz.vpshld")) {
3973 bool ZeroMask = Name[11] ==
'z';
3975 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
3976 Name.starts_with(
"avx512.mask.vpshrd") ||
3977 Name.starts_with(
"avx512.maskz.vpshrd")) {
3978 bool ZeroMask = Name[11] ==
'z';
3980 }
else if (Name ==
"sse42.crc32.64.8") {
3983 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
3985 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
3986 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
3987 Name.starts_with(
"avx512.vbroadcast.s")) {
3990 Type *EltTy = VecTy->getElementType();
3991 unsigned EltNum = VecTy->getNumElements();
3995 for (
unsigned I = 0;
I < EltNum; ++
I)
3996 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
3997 }
else if (Name.starts_with(
"sse41.pmovsx") ||
3998 Name.starts_with(
"sse41.pmovzx") ||
3999 Name.starts_with(
"avx2.pmovsx") ||
4000 Name.starts_with(
"avx2.pmovzx") ||
4001 Name.starts_with(
"avx512.mask.pmovsx") ||
4002 Name.starts_with(
"avx512.mask.pmovzx")) {
4004 unsigned NumDstElts = DstTy->getNumElements();
4008 for (
unsigned i = 0; i != NumDstElts; ++i)
4013 bool DoSext = Name.contains(
"pmovsx");
4015 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
4020 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
4021 Name ==
"avx512.mask.pmov.qd.512" ||
4022 Name ==
"avx512.mask.pmov.wb.256" ||
4023 Name ==
"avx512.mask.pmov.wb.512") {
4028 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4029 Name ==
"avx2.vbroadcasti128") {
4035 if (NumSrcElts == 2)
4038 Rep = Builder.CreateShuffleVector(
Load,
4040 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4041 Name.starts_with(
"avx512.mask.shuf.f")) {
4046 unsigned ControlBitsMask = NumLanes - 1;
4047 unsigned NumControlBits = NumLanes / 2;
4050 for (
unsigned l = 0; l != NumLanes; ++l) {
4051 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4053 if (l >= NumLanes / 2)
4054 LaneMask += NumLanes;
4055 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4056 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4062 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4063 Name.starts_with(
"avx512.mask.broadcasti")) {
4066 unsigned NumDstElts =
4070 for (
unsigned i = 0; i != NumDstElts; ++i)
4071 ShuffleMask[i] = i % NumSrcElts;
4077 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4078 Name.starts_with(
"avx2.vbroadcast") ||
4079 Name.starts_with(
"avx512.pbroadcast") ||
4080 Name.starts_with(
"avx512.mask.broadcast.s")) {
4087 Rep = Builder.CreateShuffleVector(
Op, M);
4092 }
else if (Name.starts_with(
"sse2.padds.") ||
4093 Name.starts_with(
"avx2.padds.") ||
4094 Name.starts_with(
"avx512.padds.") ||
4095 Name.starts_with(
"avx512.mask.padds.")) {
4097 }
else if (Name.starts_with(
"sse2.psubs.") ||
4098 Name.starts_with(
"avx2.psubs.") ||
4099 Name.starts_with(
"avx512.psubs.") ||
4100 Name.starts_with(
"avx512.mask.psubs.")) {
4102 }
else if (Name.starts_with(
"sse2.paddus.") ||
4103 Name.starts_with(
"avx2.paddus.") ||
4104 Name.starts_with(
"avx512.mask.paddus.")) {
4106 }
else if (Name.starts_with(
"sse2.psubus.") ||
4107 Name.starts_with(
"avx2.psubus.") ||
4108 Name.starts_with(
"avx512.mask.psubus.")) {
4110 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4115 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4119 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4124 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4129 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4130 Name ==
"avx512.psll.dq.512") {
4134 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4135 Name ==
"avx512.psrl.dq.512") {
4139 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4140 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4141 Name.starts_with(
"avx2.pblendd.")) {
4146 unsigned NumElts = VecTy->getNumElements();
4149 for (
unsigned i = 0; i != NumElts; ++i)
4150 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4152 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4153 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4154 Name ==
"avx2.vinserti128" ||
4155 Name.starts_with(
"avx512.mask.insert")) {
4159 unsigned DstNumElts =
4161 unsigned SrcNumElts =
4163 unsigned Scale = DstNumElts / SrcNumElts;
4170 for (
unsigned i = 0; i != SrcNumElts; ++i)
4172 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4173 Idxs[i] = SrcNumElts;
4174 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4188 for (
unsigned i = 0; i != DstNumElts; ++i)
4191 for (
unsigned i = 0; i != SrcNumElts; ++i)
4192 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4193 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4199 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4200 Name ==
"avx2.vextracti128" ||
4201 Name.starts_with(
"avx512.mask.vextract")) {
4204 unsigned DstNumElts =
4206 unsigned SrcNumElts =
4208 unsigned Scale = SrcNumElts / DstNumElts;
4215 for (
unsigned i = 0; i != DstNumElts; ++i) {
4216 Idxs[i] = i + (
Imm * DstNumElts);
4218 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4224 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4225 Name.starts_with(
"avx512.mask.perm.di.")) {
4229 unsigned NumElts = VecTy->getNumElements();
4232 for (
unsigned i = 0; i != NumElts; ++i)
4233 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4235 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4240 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4252 unsigned HalfSize = NumElts / 2;
4264 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4265 for (
unsigned i = 0; i < HalfSize; ++i)
4266 ShuffleMask[i] = StartIndex + i;
4269 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4270 for (
unsigned i = 0; i < HalfSize; ++i)
4271 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4273 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4275 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4276 Name.starts_with(
"avx512.mask.vpermil.p") ||
4277 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4281 unsigned NumElts = VecTy->getNumElements();
4283 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4284 unsigned IdxMask = ((1 << IdxSize) - 1);
4290 for (
unsigned i = 0; i != NumElts; ++i)
4291 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4293 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4298 }
else if (Name ==
"sse2.pshufl.w" ||
4299 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4304 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4308 for (
unsigned l = 0; l != NumElts; l += 8) {
4309 for (
unsigned i = 0; i != 4; ++i)
4310 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4311 for (
unsigned i = 4; i != 8; ++i)
4312 Idxs[i + l] = i + l;
4315 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4320 }
else if (Name ==
"sse2.pshufh.w" ||
4321 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4326 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4330 for (
unsigned l = 0; l != NumElts; l += 8) {
4331 for (
unsigned i = 0; i != 4; ++i)
4332 Idxs[i + l] = i + l;
4333 for (
unsigned i = 0; i != 4; ++i)
4334 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4337 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4342 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4349 unsigned HalfLaneElts = NumLaneElts / 2;
4352 for (
unsigned i = 0; i != NumElts; ++i) {
4354 Idxs[i] = i - (i % NumLaneElts);
4356 if ((i % NumLaneElts) >= HalfLaneElts)
4360 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4363 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4367 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4368 Name.starts_with(
"avx512.mask.movshdup") ||
4369 Name.starts_with(
"avx512.mask.movsldup")) {
4375 if (Name.starts_with(
"avx512.mask.movshdup."))
4379 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4380 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4381 Idxs[i + l + 0] = i + l +
Offset;
4382 Idxs[i + l + 1] = i + l +
Offset;
4385 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4389 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4390 Name.starts_with(
"avx512.mask.unpckl.")) {
4397 for (
int l = 0; l != NumElts; l += NumLaneElts)
4398 for (
int i = 0; i != NumLaneElts; ++i)
4399 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4401 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4405 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4406 Name.starts_with(
"avx512.mask.unpckh.")) {
4413 for (
int l = 0; l != NumElts; l += NumLaneElts)
4414 for (
int i = 0; i != NumLaneElts; ++i)
4415 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4417 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4421 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4422 Name.starts_with(
"avx512.mask.pand.")) {
4425 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4427 Rep = Builder.CreateBitCast(Rep, FTy);
4430 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4431 Name.starts_with(
"avx512.mask.pandn.")) {
4434 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4435 Rep = Builder.CreateAnd(Rep,
4437 Rep = Builder.CreateBitCast(Rep, FTy);
4440 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4441 Name.starts_with(
"avx512.mask.por.")) {
4444 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4446 Rep = Builder.CreateBitCast(Rep, FTy);
4449 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4450 Name.starts_with(
"avx512.mask.pxor.")) {
4453 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4455 Rep = Builder.CreateBitCast(Rep, FTy);
4458 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4462 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4466 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4470 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4471 if (Name.ends_with(
".512")) {
4473 if (Name[17] ==
's')
4474 IID = Intrinsic::x86_avx512_add_ps_512;
4476 IID = Intrinsic::x86_avx512_add_pd_512;
4478 Rep = Builder.CreateIntrinsic(
4486 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4487 if (Name.ends_with(
".512")) {
4489 if (Name[17] ==
's')
4490 IID = Intrinsic::x86_avx512_div_ps_512;
4492 IID = Intrinsic::x86_avx512_div_pd_512;
4494 Rep = Builder.CreateIntrinsic(
4502 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4503 if (Name.ends_with(
".512")) {
4505 if (Name[17] ==
's')
4506 IID = Intrinsic::x86_avx512_mul_ps_512;
4508 IID = Intrinsic::x86_avx512_mul_pd_512;
4510 Rep = Builder.CreateIntrinsic(
4518 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4519 if (Name.ends_with(
".512")) {
4521 if (Name[17] ==
's')
4522 IID = Intrinsic::x86_avx512_sub_ps_512;
4524 IID = Intrinsic::x86_avx512_sub_pd_512;
4526 Rep = Builder.CreateIntrinsic(
4534 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4535 Name.starts_with(
"avx512.mask.min.p")) &&
4536 Name.drop_front(18) ==
".512") {
4537 bool IsDouble = Name[17] ==
'd';
4538 bool IsMin = Name[13] ==
'i';
4540 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4541 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4544 Rep = Builder.CreateIntrinsic(
4549 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4551 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4552 {CI->getArgOperand(0), Builder.getInt1(false)});
4555 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4556 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4557 bool IsVariable = Name[16] ==
'v';
4558 char Size = Name[16] ==
'.' ? Name[17]
4559 : Name[17] ==
'.' ? Name[18]
4560 : Name[18] ==
'.' ? Name[19]
4564 if (IsVariable && Name[17] !=
'.') {
4565 if (
Size ==
'd' && Name[17] ==
'2')
4566 IID = Intrinsic::x86_avx2_psllv_q;
4567 else if (
Size ==
'd' && Name[17] ==
'4')
4568 IID = Intrinsic::x86_avx2_psllv_q_256;
4569 else if (
Size ==
's' && Name[17] ==
'4')
4570 IID = Intrinsic::x86_avx2_psllv_d;
4571 else if (
Size ==
's' && Name[17] ==
'8')
4572 IID = Intrinsic::x86_avx2_psllv_d_256;
4573 else if (
Size ==
'h' && Name[17] ==
'8')
4574 IID = Intrinsic::x86_avx512_psllv_w_128;
4575 else if (
Size ==
'h' && Name[17] ==
'1')
4576 IID = Intrinsic::x86_avx512_psllv_w_256;
4577 else if (Name[17] ==
'3' && Name[18] ==
'2')
4578 IID = Intrinsic::x86_avx512_psllv_w_512;
4581 }
else if (Name.ends_with(
".128")) {
4583 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4584 : Intrinsic::x86_sse2_psll_d;
4585 else if (
Size ==
'q')
4586 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4587 : Intrinsic::x86_sse2_psll_q;
4588 else if (
Size ==
'w')
4589 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4590 : Intrinsic::x86_sse2_psll_w;
4593 }
else if (Name.ends_with(
".256")) {
4595 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4596 : Intrinsic::x86_avx2_psll_d;
4597 else if (
Size ==
'q')
4598 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4599 : Intrinsic::x86_avx2_psll_q;
4600 else if (
Size ==
'w')
4601 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4602 : Intrinsic::x86_avx2_psll_w;
4607 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4608 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4609 : Intrinsic::x86_avx512_psll_d_512;
4610 else if (
Size ==
'q')
4611 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4612 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4613 : Intrinsic::x86_avx512_psll_q_512;
4614 else if (
Size ==
'w')
4615 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4616 : Intrinsic::x86_avx512_psll_w_512;
4622 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4623 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4624 bool IsVariable = Name[16] ==
'v';
4625 char Size = Name[16] ==
'.' ? Name[17]
4626 : Name[17] ==
'.' ? Name[18]
4627 : Name[18] ==
'.' ? Name[19]
4631 if (IsVariable && Name[17] !=
'.') {
4632 if (
Size ==
'd' && Name[17] ==
'2')
4633 IID = Intrinsic::x86_avx2_psrlv_q;
4634 else if (
Size ==
'd' && Name[17] ==
'4')
4635 IID = Intrinsic::x86_avx2_psrlv_q_256;
4636 else if (
Size ==
's' && Name[17] ==
'4')
4637 IID = Intrinsic::x86_avx2_psrlv_d;
4638 else if (
Size ==
's' && Name[17] ==
'8')
4639 IID = Intrinsic::x86_avx2_psrlv_d_256;
4640 else if (
Size ==
'h' && Name[17] ==
'8')
4641 IID = Intrinsic::x86_avx512_psrlv_w_128;
4642 else if (
Size ==
'h' && Name[17] ==
'1')
4643 IID = Intrinsic::x86_avx512_psrlv_w_256;
4644 else if (Name[17] ==
'3' && Name[18] ==
'2')
4645 IID = Intrinsic::x86_avx512_psrlv_w_512;
4648 }
else if (Name.ends_with(
".128")) {
4650 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4651 : Intrinsic::x86_sse2_psrl_d;
4652 else if (
Size ==
'q')
4653 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4654 : Intrinsic::x86_sse2_psrl_q;
4655 else if (
Size ==
'w')
4656 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4657 : Intrinsic::x86_sse2_psrl_w;
4660 }
else if (Name.ends_with(
".256")) {
4662 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4663 : Intrinsic::x86_avx2_psrl_d;
4664 else if (
Size ==
'q')
4665 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4666 : Intrinsic::x86_avx2_psrl_q;
4667 else if (
Size ==
'w')
4668 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4669 : Intrinsic::x86_avx2_psrl_w;
4674 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4675 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4676 : Intrinsic::x86_avx512_psrl_d_512;
4677 else if (
Size ==
'q')
4678 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4679 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4680 : Intrinsic::x86_avx512_psrl_q_512;
4681 else if (
Size ==
'w')
4682 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4683 : Intrinsic::x86_avx512_psrl_w_512;
4689 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4690 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4691 bool IsVariable = Name[16] ==
'v';
4692 char Size = Name[16] ==
'.' ? Name[17]
4693 : Name[17] ==
'.' ? Name[18]
4694 : Name[18] ==
'.' ? Name[19]
4698 if (IsVariable && Name[17] !=
'.') {
4699 if (
Size ==
's' && Name[17] ==
'4')
4700 IID = Intrinsic::x86_avx2_psrav_d;
4701 else if (
Size ==
's' && Name[17] ==
'8')
4702 IID = Intrinsic::x86_avx2_psrav_d_256;
4703 else if (
Size ==
'h' && Name[17] ==
'8')
4704 IID = Intrinsic::x86_avx512_psrav_w_128;
4705 else if (
Size ==
'h' && Name[17] ==
'1')
4706 IID = Intrinsic::x86_avx512_psrav_w_256;
4707 else if (Name[17] ==
'3' && Name[18] ==
'2')
4708 IID = Intrinsic::x86_avx512_psrav_w_512;
4711 }
else if (Name.ends_with(
".128")) {
4713 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4714 : Intrinsic::x86_sse2_psra_d;
4715 else if (
Size ==
'q')
4716 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4717 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4718 : Intrinsic::x86_avx512_psra_q_128;
4719 else if (
Size ==
'w')
4720 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4721 : Intrinsic::x86_sse2_psra_w;
4724 }
else if (Name.ends_with(
".256")) {
4726 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4727 : Intrinsic::x86_avx2_psra_d;
4728 else if (
Size ==
'q')
4729 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4730 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4731 : Intrinsic::x86_avx512_psra_q_256;
4732 else if (
Size ==
'w')
4733 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4734 : Intrinsic::x86_avx2_psra_w;
4739 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4740 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4741 : Intrinsic::x86_avx512_psra_d_512;
4742 else if (
Size ==
'q')
4743 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4744 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4745 : Intrinsic::x86_avx512_psra_q_512;
4746 else if (
Size ==
'w')
4747 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4748 : Intrinsic::x86_avx512_psra_w_512;
4754 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4756 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4758 }
else if (Name.ends_with(
".movntdqa")) {
4762 LoadInst *LI = Builder.CreateAlignedLoad(
4767 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4768 Name.starts_with(
"fma.vfmsub.") ||
4769 Name.starts_with(
"fma.vfnmadd.") ||
4770 Name.starts_with(
"fma.vfnmsub.")) {
4771 bool NegMul = Name[6] ==
'n';
4772 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4773 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4784 if (NegMul && !IsScalar)
4785 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4786 if (NegMul && IsScalar)
4787 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4789 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4791 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4795 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4803 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4807 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4808 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4809 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4810 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4811 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4812 bool IsMask3 = Name[11] ==
'3';
4813 bool IsMaskZ = Name[11] ==
'z';
4815 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4816 bool NegMul = Name[2] ==
'n';
4817 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4823 if (NegMul && (IsMask3 || IsMaskZ))
4824 A = Builder.CreateFNeg(
A);
4825 if (NegMul && !(IsMask3 || IsMaskZ))
4826 B = Builder.CreateFNeg(
B);
4828 C = Builder.CreateFNeg(
C);
4830 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4831 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4832 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4839 if (Name.back() ==
'd')
4840 IID = Intrinsic::x86_avx512_vfmadd_f64;
4842 IID = Intrinsic::x86_avx512_vfmadd_f32;
4843 Rep = Builder.CreateIntrinsic(IID,
Ops);
4845 Rep = Builder.CreateFMA(
A,
B,
C);
4854 if (NegAcc && IsMask3)
4859 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4861 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4862 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4863 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4864 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4865 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4866 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4867 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4868 bool IsMask3 = Name[11] ==
'3';
4869 bool IsMaskZ = Name[11] ==
'z';
4871 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4872 bool NegMul = Name[2] ==
'n';
4873 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4879 if (NegMul && (IsMask3 || IsMaskZ))
4880 A = Builder.CreateFNeg(
A);
4881 if (NegMul && !(IsMask3 || IsMaskZ))
4882 B = Builder.CreateFNeg(
B);
4884 C = Builder.CreateFNeg(
C);
4891 if (Name[Name.size() - 5] ==
's')
4892 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4894 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4898 Rep = Builder.CreateFMA(
A,
B,
C);
4906 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4910 if (VecWidth == 128 && EltWidth == 32)
4911 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4912 else if (VecWidth == 256 && EltWidth == 32)
4913 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4914 else if (VecWidth == 128 && EltWidth == 64)
4915 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4916 else if (VecWidth == 256 && EltWidth == 64)
4917 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4923 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4924 Rep = Builder.CreateIntrinsic(IID,
Ops);
4925 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4926 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4927 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4928 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4929 bool IsMask3 = Name[11] ==
'3';
4930 bool IsMaskZ = Name[11] ==
'z';
4932 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4933 bool IsSubAdd = Name[3] ==
's';
4937 if (Name[Name.size() - 5] ==
's')
4938 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4940 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4945 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4947 Rep = Builder.CreateIntrinsic(IID,
Ops);
4956 Value *Odd = Builder.CreateCall(FMA,
Ops);
4957 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4958 Value *Even = Builder.CreateCall(FMA,
Ops);
4964 for (
int i = 0; i != NumElts; ++i)
4965 Idxs[i] = i + (i % 2) * NumElts;
4967 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
4975 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
4976 Name.starts_with(
"avx512.maskz.pternlog.")) {
4977 bool ZeroMask = Name[11] ==
'z';
4981 if (VecWidth == 128 && EltWidth == 32)
4982 IID = Intrinsic::x86_avx512_pternlog_d_128;
4983 else if (VecWidth == 256 && EltWidth == 32)
4984 IID = Intrinsic::x86_avx512_pternlog_d_256;
4985 else if (VecWidth == 512 && EltWidth == 32)
4986 IID = Intrinsic::x86_avx512_pternlog_d_512;
4987 else if (VecWidth == 128 && EltWidth == 64)
4988 IID = Intrinsic::x86_avx512_pternlog_q_128;
4989 else if (VecWidth == 256 && EltWidth == 64)
4990 IID = Intrinsic::x86_avx512_pternlog_q_256;
4991 else if (VecWidth == 512 && EltWidth == 64)
4992 IID = Intrinsic::x86_avx512_pternlog_q_512;
4998 Rep = Builder.CreateIntrinsic(IID, Args);
5002 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
5003 Name.starts_with(
"avx512.maskz.vpmadd52")) {
5004 bool ZeroMask = Name[11] ==
'z';
5005 bool High = Name[20] ==
'h' || Name[21] ==
'h';
5008 if (VecWidth == 128 && !
High)
5009 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
5010 else if (VecWidth == 256 && !
High)
5011 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
5012 else if (VecWidth == 512 && !
High)
5013 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
5014 else if (VecWidth == 128 &&
High)
5015 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
5016 else if (VecWidth == 256 &&
High)
5017 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
5018 else if (VecWidth == 512 &&
High)
5019 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5025 Rep = Builder.CreateIntrinsic(IID, Args);
5029 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5030 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5031 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5032 bool ZeroMask = Name[11] ==
'z';
5033 bool IndexForm = Name[17] ==
'i';
5035 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5036 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5037 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5038 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5039 bool ZeroMask = Name[11] ==
'z';
5040 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5043 if (VecWidth == 128 && !IsSaturating)
5044 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5045 else if (VecWidth == 256 && !IsSaturating)
5046 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5047 else if (VecWidth == 512 && !IsSaturating)
5048 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5049 else if (VecWidth == 128 && IsSaturating)
5050 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5051 else if (VecWidth == 256 && IsSaturating)
5052 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5053 else if (VecWidth == 512 && IsSaturating)
5054 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5064 if (Args[1]->
getType()->isVectorTy() &&
5067 ->isIntegerTy(32) &&
5068 Args[2]->
getType()->isVectorTy() &&
5071 ->isIntegerTy(32)) {
5072 Type *NewArgType =
nullptr;
5073 if (VecWidth == 128)
5075 else if (VecWidth == 256)
5077 else if (VecWidth == 512)
5083 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5084 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5087 Rep = Builder.CreateIntrinsic(IID, Args);
5091 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5092 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5093 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5094 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5095 bool ZeroMask = Name[11] ==
'z';
5096 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5099 if (VecWidth == 128 && !IsSaturating)
5100 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5101 else if (VecWidth == 256 && !IsSaturating)
5102 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5103 else if (VecWidth == 512 && !IsSaturating)
5104 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5105 else if (VecWidth == 128 && IsSaturating)
5106 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5107 else if (VecWidth == 256 && IsSaturating)
5108 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5109 else if (VecWidth == 512 && IsSaturating)
5110 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5120 if (Args[1]->
getType()->isVectorTy() &&
5123 ->isIntegerTy(32) &&
5124 Args[2]->
getType()->isVectorTy() &&
5127 ->isIntegerTy(32)) {
5128 Type *NewArgType =
nullptr;
5129 if (VecWidth == 128)
5131 else if (VecWidth == 256)
5133 else if (VecWidth == 512)
5139 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5140 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5143 Rep = Builder.CreateIntrinsic(IID, Args);
5147 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5148 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5149 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5151 if (Name[0] ==
'a' && Name.back() ==
'2')
5152 IID = Intrinsic::x86_addcarry_32;
5153 else if (Name[0] ==
'a' && Name.back() ==
'4')
5154 IID = Intrinsic::x86_addcarry_64;
5155 else if (Name[0] ==
's' && Name.back() ==
'2')
5156 IID = Intrinsic::x86_subborrow_32;
5157 else if (Name[0] ==
's' && Name.back() ==
'4')
5158 IID = Intrinsic::x86_subborrow_64;
5165 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5168 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5171 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5175 }
else if (Name.starts_with(
"avx512.mask.") &&
5178 }
else if (Name.starts_with(
"bmi.pdep.")) {
5180 }
else if (Name.starts_with(
"bmi.pext.")) {
5190 if (Name.starts_with(
"neon.bfcvt")) {
5191 if (Name.starts_with(
"neon.bfcvtn2")) {
5193 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5195 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5196 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5199 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5200 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5202 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5206 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5207 return Builder.CreateShuffleVector(
5210 return Builder.CreateFPTrunc(CI->
getOperand(0),
5213 }
else if (Name.starts_with(
"sve.fcvt")) {
5216 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5217 .
Case(
"sve.fcvtnt.bf16f32",
5218 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5230 if (Args[1]->
getType() != BadPredTy)
5233 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5234 BadPredTy, Args[1]);
5235 Args[1] = Builder.CreateIntrinsic(
5236 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5238 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5242 if (Name ==
"neon.vcvtfp2hf")
5243 return Builder.CreateBitCast(
5244 Builder.CreateFPTrunc(
5248 if (Name ==
"neon.vcvthf2fp")
5249 return Builder.CreateFPExt(
5250 Builder.CreateBitCast(
5260 if (Name ==
"mve.vctp64.old") {
5263 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5266 Value *C1 = Builder.CreateIntrinsic(
5267 Intrinsic::arm_mve_pred_v2i,
5269 return Builder.CreateIntrinsic(
5270 Intrinsic::arm_mve_pred_i2v,
5272 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5273 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5274 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5275 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5277 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5278 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5279 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5280 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5282 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5283 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5284 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5285 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5286 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5287 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5288 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5289 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5290 std::vector<Type *> Tys;
5294 case Intrinsic::arm_mve_mull_int_predicated:
5295 case Intrinsic::arm_mve_vqdmull_predicated:
5296 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5299 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5300 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5301 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5305 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5309 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5313 case Intrinsic::arm_cde_vcx1q_predicated:
5314 case Intrinsic::arm_cde_vcx1qa_predicated:
5315 case Intrinsic::arm_cde_vcx2q_predicated:
5316 case Intrinsic::arm_cde_vcx2qa_predicated:
5317 case Intrinsic::arm_cde_vcx3q_predicated:
5318 case Intrinsic::arm_cde_vcx3qa_predicated:
5325 std::vector<Value *>
Ops;
5327 Type *Ty =
Op->getType();
5328 if (Ty->getScalarSizeInBits() == 1) {
5329 Value *C1 = Builder.CreateIntrinsic(
5330 Intrinsic::arm_mve_pred_v2i,
5332 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty}, C1);
5337 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5352 auto UpgradeLegacyWMMAIUIntrinsicCall =
5357 Args.push_back(Builder.getFalse());
5361 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5368 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5372 NewCall->copyMetadata(*CI);
5376 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5377 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5378 "intrinsic should have 7 arguments");
5381 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5383 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5384 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5385 "intrinsic should have 8 arguments");
5390 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5393 switch (
F->getIntrinsicID()) {
5396 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5397 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5398 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5399 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5400 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5401 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5416 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5419 F->getParent(),
F->getIntrinsicID(), Overloads);
5424 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5428 NewCall->copyMetadata(*CI);
5429 NewCall->takeName(CI);
5434 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5440 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5441 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5449 if (Name.starts_with(
"addrspacecast.nonnull")) {
5452 Value *ASC = Builder.CreateAddrSpaceCast(
5475 if (NumOperands < 3)
5488 bool IsVolatile =
false;
5492 if (NumOperands > 3)
5497 if (NumOperands > 5) {
5499 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5513 if (VT->getElementType()->isIntegerTy(16)) {
5516 Val = Builder.CreateBitCast(Val, AsBF16);
5524 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5526 unsigned AddrSpace = PtrTy->getAddressSpace();
5529 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5531 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5536 MDNode *RangeNotPrivate =
5539 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5545 return Builder.CreateBitCast(RMW, RetTy);
5566 return MAV->getMetadata();
5575 if (Name ==
"label") {
5577 }
else if (Name ==
"assign") {
5584 }
else if (Name ==
"declare") {
5588 }
else if (Name ==
"addr") {
5598 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5599 }
else if (Name ==
"value") {
5602 unsigned ExprOp = 2;
5617 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5625 int64_t OffsetVal =
Offset->getSExtValue();
5626 return Builder.CreateIntrinsic(OffsetVal >= 0
5627 ? Intrinsic::vector_splice_left
5628 : Intrinsic::vector_splice_right,
5630 {CI->getArgOperand(0), CI->getArgOperand(1),
5631 Builder.getInt32(std::abs(OffsetVal))});
5636 if (Name.starts_with(
"to.fp16")) {
5638 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5639 return Builder.CreateBitCast(Cast, CI->
getType());
5642 if (Name.starts_with(
"from.fp16")) {
5644 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5645 return Builder.CreateFPExt(Cast, CI->
getType());
5704 else if (Opcode == Instruction::ICmp)
5707 else if (Opcode == Instruction::FCmp)
5710 else if (Opcode == Instruction::Select)
5715 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5727 if (Defaults.empty())
5730 unsigned OldArgCount = CI->
arg_size();
5731 unsigned NewArgCount = NewFn->
arg_size();
5733 if (OldArgCount < FirstDefault)
5737 if (OldArgCount > NewArgCount)
5742 if (OldArgCount == NewArgCount) {
5754 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5755 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5756 "missing argument outside the default range");
5757 Type *ParamTy = NewFT->getParamType(Idx);
5762 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5768 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5799 if (!Name.consume_front(
"llvm."))
5802 bool IsX86 = Name.consume_front(
"x86.");
5803 bool IsNVVM = Name.consume_front(
"nvvm.");
5804 bool IsAArch64 = Name.consume_front(
"aarch64.");
5805 bool IsARM = Name.consume_front(
"arm.");
5806 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5807 bool IsDbg = Name.consume_front(
"dbg.");
5809 (Name.consume_front(
"experimental.vector.splice") ||
5810 Name.consume_front(
"vector.splice")) &&
5811 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5812 Value *Rep =
nullptr;
5814 if (!IsX86 && Name ==
"stackprotectorcheck") {
5816 }
else if (IsNVVM) {
5820 }
else if (IsAArch64) {
5824 }
else if (IsAMDGCN) {
5828 }
else if (IsOldSplice) {
5830 }
else if (Name.consume_front(
"convert.")) {
5832 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5847 const auto &DefaultCase = [&]() ->
void {
5855 "Unknown function for CallBase upgrade and isn't just a name change");
5863 "Return type must have changed");
5864 assert(OldST->getNumElements() ==
5866 "Must have same number of elements");
5869 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5872 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5873 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5874 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5895 case Intrinsic::arm_neon_vst1:
5896 case Intrinsic::arm_neon_vst2:
5897 case Intrinsic::arm_neon_vst3:
5898 case Intrinsic::arm_neon_vst4:
5899 case Intrinsic::arm_neon_vst2lane:
5900 case Intrinsic::arm_neon_vst3lane:
5901 case Intrinsic::arm_neon_vst4lane: {
5903 NewCall = Builder.CreateCall(NewFn, Args);
5906 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5907 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5908 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5913 NewCall = Builder.CreateCall(NewFn, Args);
5916 case Intrinsic::aarch64_sve_ld3_sret:
5917 case Intrinsic::aarch64_sve_ld4_sret:
5918 case Intrinsic::aarch64_sve_ld2_sret: {
5926 Name = Name.substr(5);
5933 unsigned MinElts = RetTy->getMinNumElements() /
N;
5935 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5937 for (
unsigned I = 0;
I <
N;
I++) {
5938 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5939 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5945 case Intrinsic::coro_end_async:
5946 case Intrinsic::coro_end: {
5948 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5950 NewCall = Builder.CreateCall(NewFn, Args);
5955 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5956 Value *InRamp = Builder.CreateCall(IsInRamp);
5966 case Intrinsic::vector_extract: {
5968 Name = Name.substr(5);
5969 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
5974 unsigned MinElts = RetTy->getMinNumElements();
5977 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
5981 case Intrinsic::vector_insert: {
5983 Name = Name.substr(5);
5984 if (!Name.starts_with(
"aarch64.sve.tuple")) {
5988 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
5993 NewCall = Builder.CreateCall(
5997 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
6003 assert(
N > 1 &&
"Create is expected to be between 2-4");
6006 unsigned MinElts = RetTy->getMinNumElements() /
N;
6007 for (
unsigned I = 0;
I <
N;
I++) {
6009 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
6016 case Intrinsic::arm_neon_bfdot:
6017 case Intrinsic::arm_neon_bfmmla:
6018 case Intrinsic::arm_neon_bfmlalb:
6019 case Intrinsic::arm_neon_bfmlalt:
6020 case Intrinsic::aarch64_neon_bfdot:
6021 case Intrinsic::aarch64_neon_bfmmla:
6022 case Intrinsic::aarch64_neon_bfmlalb:
6023 case Intrinsic::aarch64_neon_bfmlalt: {
6026 "Mismatch between function args and call args");
6027 size_t OperandWidth =
6029 assert((OperandWidth == 64 || OperandWidth == 128) &&
6030 "Unexpected operand width");
6032 auto Iter = CI->
args().begin();
6033 Args.push_back(*Iter++);
6034 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6035 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6036 NewCall = Builder.CreateCall(NewFn, Args);
6040 case Intrinsic::bitreverse:
6041 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6044 case Intrinsic::ctlz:
6045 case Intrinsic::cttz: {
6052 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6056 case Intrinsic::objectsize: {
6057 Value *NullIsUnknownSize =
6061 NewCall = Builder.CreateCall(
6066 case Intrinsic::ctpop:
6067 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6069 case Intrinsic::dbg_value: {
6071 Name = Name.substr(5);
6073 if (Name.starts_with(
"dbg.addr")) {
6087 if (
Offset->isNullValue()) {
6088 NewCall = Builder.CreateCall(
6097 case Intrinsic::ptr_annotation:
6105 NewCall = Builder.CreateCall(
6114 case Intrinsic::var_annotation:
6121 NewCall = Builder.CreateCall(
6130 case Intrinsic::riscv_aes32dsi:
6131 case Intrinsic::riscv_aes32dsmi:
6132 case Intrinsic::riscv_aes32esi:
6133 case Intrinsic::riscv_aes32esmi:
6134 case Intrinsic::riscv_sm4ks:
6135 case Intrinsic::riscv_sm4ed: {
6145 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6146 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6152 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6153 Value *Res = NewCall;
6155 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6161 case Intrinsic::nvvm_mapa_shared_cluster: {
6165 Value *Res = NewCall;
6166 Res = Builder.CreateAddrSpaceCast(
6173 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6175 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6177 Args[0] = Builder.CreateAddrSpaceCast(
6181 Args.push_back(Builder.getInt32(0));
6183 NewCall = Builder.CreateCall(NewFn, Args);
6189 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6194 for (
unsigned I = 0;
I < 4; ++
I)
6196 Args.push_back(Builder.getInt32(0));
6197 Args.push_back(Builder.getInt32(0));
6200 Args.push_back(Builder.getInt1(
false));
6201 Args.push_back(Builder.getInt32(0));
6203 NewCall = Builder.CreateCall(NewFn, Args);
6209 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6212 Args[0] = Builder.CreateAddrSpaceCast(
6215 NewCall = Builder.CreateCall(NewFn, Args);
6222#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6223 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6225#undef G2S_CLUSTER_CASE
6230 Args[0] = Builder.CreateAddrSpaceCast(
6236 Args.push_back(Builder.getInt32(0));
6238 NewCall = Builder.CreateCall(NewFn, Args);
6245#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6246 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6254 "expected only the trailing flag_valid_pattern to be missing");
6255 Args.push_back(Builder.getInt32(0));
6257 NewCall = Builder.CreateCall(NewFn, Args);
6263#undef NVVM_TMA_G2S_MODES
6266 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6267 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6268 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6269 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6270 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6271 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6272 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6273 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6275 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6279 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6280 NewCall = Builder.CreateCall(NewFn, Args);
6283 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6284 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6285 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6286 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6289 Builder.getFalse()});
6291 case Intrinsic::nvvm_mbarrier_init: {
6295 if (Args.size() == 2)
6296 Args.push_back(Builder.getInt32(0));
6297 NewCall = Builder.CreateCall(NewFn, Args);
6300 case Intrinsic::riscv_sha256sig0:
6301 case Intrinsic::riscv_sha256sig1:
6302 case Intrinsic::riscv_sha256sum0:
6303 case Intrinsic::riscv_sha256sum1:
6304 case Intrinsic::riscv_sm3p0:
6305 case Intrinsic::riscv_sm3p1: {
6312 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6314 NewCall = Builder.CreateCall(NewFn, Arg);
6316 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6323 case Intrinsic::x86_xop_vfrcz_ss:
6324 case Intrinsic::x86_xop_vfrcz_sd:
6325 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6328 case Intrinsic::x86_xop_vpermil2pd:
6329 case Intrinsic::x86_xop_vpermil2ps:
6330 case Intrinsic::x86_xop_vpermil2pd_256:
6331 case Intrinsic::x86_xop_vpermil2ps_256: {
6335 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6336 NewCall = Builder.CreateCall(NewFn, Args);
6340 case Intrinsic::x86_sse41_ptestc:
6341 case Intrinsic::x86_sse41_ptestz:
6342 case Intrinsic::x86_sse41_ptestnzc: {
6356 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6357 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6359 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6363 case Intrinsic::x86_rdtscp: {
6369 NewCall = Builder.CreateCall(NewFn);
6371 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6374 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6382 case Intrinsic::x86_sse41_insertps:
6383 case Intrinsic::x86_sse41_dppd:
6384 case Intrinsic::x86_sse41_dpps:
6385 case Intrinsic::x86_sse41_mpsadbw:
6386 case Intrinsic::x86_avx_dp_ps_256:
6387 case Intrinsic::x86_avx2_mpsadbw: {
6393 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6394 NewCall = Builder.CreateCall(NewFn, Args);
6398 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6399 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6400 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6401 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6402 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6403 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6409 NewCall = Builder.CreateCall(NewFn, Args);
6418 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6419 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6420 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6421 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6422 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6423 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6427 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6428 Args[1] = Builder.CreateBitCast(
6431 NewCall = Builder.CreateCall(NewFn, Args);
6432 Value *Res = Builder.CreateBitCast(
6440 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6441 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6442 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6446 Args[1] = Builder.CreateBitCast(
6448 Args[2] = Builder.CreateBitCast(
6451 NewCall = Builder.CreateCall(NewFn, Args);
6455 case Intrinsic::thread_pointer: {
6456 NewCall = Builder.CreateCall(NewFn, {});
6460 case Intrinsic::memcpy:
6461 case Intrinsic::memmove:
6462 case Intrinsic::memset: {
6478 NewCall = Builder.CreateCall(NewFn, Args);
6481 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6482 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6483 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6488 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6491 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6495 case Intrinsic::masked_load:
6496 case Intrinsic::masked_gather:
6497 case Intrinsic::masked_store:
6498 case Intrinsic::masked_scatter: {
6504 auto GetMaybeAlign = [](
Value *
Op) {
6506 uint64_t Val = CI->getZExtValue();
6514 auto GetAlign = [&](
Value *
Op) {
6523 case Intrinsic::masked_load:
6524 NewCall = Builder.CreateMaskedLoad(
6528 case Intrinsic::masked_gather:
6529 NewCall = Builder.CreateMaskedGather(
6535 case Intrinsic::masked_store:
6536 NewCall = Builder.CreateMaskedStore(
6540 case Intrinsic::masked_scatter:
6541 NewCall = Builder.CreateMaskedScatter(
6543 DL.getValueOrABITypeAlignment(
6557 case Intrinsic::lifetime_start:
6558 case Intrinsic::lifetime_end: {
6570 NewCall = Builder.CreateLifetimeStart(Ptr);
6572 NewCall = Builder.CreateLifetimeEnd(Ptr);
6581 case Intrinsic::x86_avx512_vpdpbusd_128:
6582 case Intrinsic::x86_avx512_vpdpbusd_256:
6583 case Intrinsic::x86_avx512_vpdpbusd_512:
6584 case Intrinsic::x86_avx512_vpdpbusds_128:
6585 case Intrinsic::x86_avx512_vpdpbusds_256:
6586 case Intrinsic::x86_avx512_vpdpbusds_512:
6587 case Intrinsic::x86_avx2_vpdpbssd_128:
6588 case Intrinsic::x86_avx2_vpdpbssd_256:
6589 case Intrinsic::x86_avx10_vpdpbssd_512:
6590 case Intrinsic::x86_avx2_vpdpbssds_128:
6591 case Intrinsic::x86_avx2_vpdpbssds_256:
6592 case Intrinsic::x86_avx10_vpdpbssds_512:
6593 case Intrinsic::x86_avx2_vpdpbsud_128:
6594 case Intrinsic::x86_avx2_vpdpbsud_256:
6595 case Intrinsic::x86_avx10_vpdpbsud_512:
6596 case Intrinsic::x86_avx2_vpdpbsuds_128:
6597 case Intrinsic::x86_avx2_vpdpbsuds_256:
6598 case Intrinsic::x86_avx10_vpdpbsuds_512:
6599 case Intrinsic::x86_avx2_vpdpbuud_128:
6600 case Intrinsic::x86_avx2_vpdpbuud_256:
6601 case Intrinsic::x86_avx10_vpdpbuud_512:
6602 case Intrinsic::x86_avx2_vpdpbuuds_128:
6603 case Intrinsic::x86_avx2_vpdpbuuds_256:
6604 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6609 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6610 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6612 NewCall = Builder.CreateCall(NewFn, Args);
6615 case Intrinsic::x86_avx512_vpdpwssd_128:
6616 case Intrinsic::x86_avx512_vpdpwssd_256:
6617 case Intrinsic::x86_avx512_vpdpwssd_512:
6618 case Intrinsic::x86_avx512_vpdpwssds_128:
6619 case Intrinsic::x86_avx512_vpdpwssds_256:
6620 case Intrinsic::x86_avx512_vpdpwssds_512:
6621 case Intrinsic::x86_avx2_vpdpwsud_128:
6622 case Intrinsic::x86_avx2_vpdpwsud_256:
6623 case Intrinsic::x86_avx10_vpdpwsud_512:
6624 case Intrinsic::x86_avx2_vpdpwsuds_128:
6625 case Intrinsic::x86_avx2_vpdpwsuds_256:
6626 case Intrinsic::x86_avx10_vpdpwsuds_512:
6627 case Intrinsic::x86_avx2_vpdpwusd_128:
6628 case Intrinsic::x86_avx2_vpdpwusd_256:
6629 case Intrinsic::x86_avx10_vpdpwusd_512:
6630 case Intrinsic::x86_avx2_vpdpwusds_128:
6631 case Intrinsic::x86_avx2_vpdpwusds_256:
6632 case Intrinsic::x86_avx10_vpdpwusds_512:
6633 case Intrinsic::x86_avx2_vpdpwuud_128:
6634 case Intrinsic::x86_avx2_vpdpwuud_256:
6635 case Intrinsic::x86_avx10_vpdpwuud_512:
6636 case Intrinsic::x86_avx2_vpdpwuuds_128:
6637 case Intrinsic::x86_avx2_vpdpwuuds_256:
6638 case Intrinsic::x86_avx10_vpdpwuuds_512:
6643 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6644 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6646 NewCall = Builder.CreateCall(NewFn, Args);
6649 assert(NewCall &&
"Should have either set this variable or returned through "
6650 "the default case");
6657 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6671 F->eraseFromParent();
6677 if (NumOperands == 0)
6685 if (NumOperands == 3) {
6689 Metadata *Elts2[] = {ScalarType, ScalarType,
6705 if (NumOperands == 0 || NumOperands % 3 != 0)
6710 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6715 if (Upgraded ==
Tag)
6725 if (
Opc != Instruction::BitCast)
6729 Type *SrcTy = V->getType();
6746 if (
Opc != Instruction::BitCast)
6749 Type *SrcTy =
C->getType();
6766 if (Flag.getNumOperands() < 3)
6767 return std::nullopt;
6769 return Name->getString();
6770 return std::nullopt;
6776 if (!
getIROptions(M.getContext()).auto_upgrade_debug_info)
6784 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6785 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6786 if (auto Name = getModuleFlagNameSafely(*Flag))
6787 return *Name ==
"Debug Info Version";
6790 if (OpIt != ModFlags->op_end()) {
6791 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6798 bool BrokenDebugInfo =
false;
6801 if (!BrokenDebugInfo)
6807 M.getContext().diagnose(Diag);
6814 M.getContext().diagnose(DiagVersion);
6824 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6827 if (
F->hasFnAttribute(Attr)) {
6830 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6832 auto [Part, Rest] = S.
split(
',');
6838 const unsigned Dim = DimC -
'x';
6839 assert(Dim < 3 &&
"Unexpected dim char");
6849 F->addFnAttr(Attr, NewAttr);
6853 return S ==
"x" || S ==
"y" || S ==
"z";
6858 if (
K ==
"kernel") {
6870 const unsigned Idx = (AlignIdxValuePair >> 16);
6871 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6876 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6881 if (
K ==
"minctasm") {
6886 if (
K ==
"maxnreg") {
6891 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6895 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6899 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6903 if (
K ==
"grid_constant") {
6918 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6925 if (!SeenNodes.
insert(MD).second)
6932 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6939 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6941 const MDOperand &V = MD->getOperand(j + 1);
6947 if (NewOperands.
size() > 1)
6960 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
6961 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
6962 if (ModRetainReleaseMarker) {
6968 ID->getString().split(ValueComp,
"#");
6969 if (ValueComp.
size() == 2) {
6970 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
6974 M.eraseNamedMetadata(ModRetainReleaseMarker);
6985 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
7011 bool InvalidCast =
false;
7013 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7026 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7028 Args.push_back(Arg);
7035 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7040 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7053 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7061 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7062 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7063 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7064 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7065 {
"objc_autoreleaseReturnValue",
7066 llvm::Intrinsic::objc_autoreleaseReturnValue},
7067 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7068 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7069 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7070 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7071 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7072 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7073 {
"objc_release", llvm::Intrinsic::objc_release},
7074 {
"objc_retain", llvm::Intrinsic::objc_retain},
7075 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7076 {
"objc_retainAutoreleaseReturnValue",
7077 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7078 {
"objc_retainAutoreleasedReturnValue",
7079 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7080 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7081 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7082 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7083 {
"objc_unsafeClaimAutoreleasedReturnValue",
7084 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7085 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7086 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7087 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7088 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7089 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7090 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7091 {
"objc_arc_annotation_topdown_bbstart",
7092 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7093 {
"objc_arc_annotation_topdown_bbend",
7094 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7095 {
"objc_arc_annotation_bottomup_bbstart",
7096 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7097 {
"objc_arc_annotation_bottomup_bbend",
7098 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7100 for (
auto &
I : RuntimeFuncs)
7101 UpgradeToIntrinsic(
I.first,
I.second);
7125 std::optional<bool> UseAddressDisc;
7128 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7129 for (
const MDNode *Flag : ModFlags->operands()) {
7131 if (Name && (*Name ==
"ptrauth-init-fini" ||
7132 *Name ==
"ptrauth-init-fini-address-discrimination"))
7137 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7138 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7139 constexpr unsigned ExpectedAddressMarker = 1;
7142 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7145 bool HasAddressDisc;
7146 if (!CPA->hasAddressDiscriminator())
7147 HasAddressDisc =
false;
7148 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7149 HasAddressDisc =
true;
7153 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7156 UseAddressDisc = HasAddressDisc;
7157 return CPA->getPointer();
7161 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7164 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7166 if (!GV || !GV->hasInitializer())
7170 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7173 std::vector<Constant *> NewStructors;
7174 NewStructors.reserve(OldStructorsArray->getNumOperands());
7176 for (
Use &U : OldStructorsArray->operands()) {
7185 Func = UpgradeSinglePointer(Func);
7189 NewStructors.push_back(
7198 if (GlobalArraysToUpgrade.
empty())
7200 assert(UseAddressDisc.has_value());
7202 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7203 GV->setInitializer(NewInit);
7206 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7216 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7220 bool HasObjCFlag =
false, HasClassProperties =
false;
7221 bool HasSwiftVersionFlag =
false;
7222 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7229 if (
Op->getNumOperands() != 3)
7243 if (ID->getString() ==
"Objective-C Image Info Version")
7245 if (ID->getString() ==
"Objective-C Class Properties")
7246 HasClassProperties =
true;
7248 if (ID->getString() ==
"PIC Level") {
7249 if (
auto *Behavior =
7251 uint64_t V = Behavior->getLimitedValue();
7257 if (ID->getString() ==
"PIE Level")
7258 if (
auto *Behavior =
7266 if (ID->getString() ==
"branch-target-enforcement" ||
7267 (ID->getString().starts_with(
"sign-return-address") &&
7268 ID->getString() !=
"sign-return-address-harden")) {
7269 if (
auto *Behavior =
7275 Op->getOperand(1),
Op->getOperand(2)};
7285 if (ID->getString() ==
"Objective-C Image Info Section") {
7288 Value->getString().split(ValueComp,
" ");
7289 if (ValueComp.
size() != 1) {
7290 std::string NewValue;
7291 for (
auto &S : ValueComp)
7292 NewValue += S.str();
7303 if (ID->getString() ==
"Objective-C Garbage Collection") {
7306 assert(Md->getValue() &&
"Expected non-empty metadata");
7307 auto Type = Md->getValue()->getType();
7310 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7311 if ((Val & 0xff) != Val) {
7312 HasSwiftVersionFlag =
true;
7313 SwiftABIVersion = (Val & 0xff00) >> 8;
7314 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7315 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7326 if (ID->getString() ==
"amdgpu_code_object_version") {
7329 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7338 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7367 if (HasObjCFlag && !HasClassProperties) {
7373 if (HasSwiftVersionFlag) {
7377 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7379 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7387 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7391 auto MatchesVersion = [](
const MDNode *
Op) {
7392 return Op->getNumOperands() >= 3 &&
7406 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7407 assert(
Op->getNumOperands() >= 2 &&
7408 "Expected at least 2 operands - name and linkage type");
7420 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7431 auto TrimSpaces = [](
StringRef Section) -> std::string {
7433 Section.split(Components,
',');
7438 for (
auto Component : Components)
7439 OS <<
',' << Component.trim();
7444 for (
auto &GV : M.globals()) {
7445 if (!GV.hasSection())
7450 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7455 GV.setSection(TrimSpaces(Section));
7471struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7472 StrictFPUpgradeVisitor() =
default;
7475 if (!
Call.isStrictFP())
7481 Call.removeFnAttr(Attribute::StrictFP);
7482 Call.addFnAttr(Attribute::NoBuiltin);
7487struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7488 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7489 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7491 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7506 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7507 StrictFPUpgradeVisitor SFPV;
7512 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7513 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7514 for (
auto &Arg :
F.args())
7516 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7518 bool AddingAttrs =
false, RemovingAttrs =
false;
7519 AttrBuilder AttrsToAdd(
F.getContext());
7524 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7525 A.isValid() &&
A.isStringAttribute()) {
7526 F.setSection(
A.getValueAsString());
7528 RemovingAttrs =
true;
7532 A.isValid() &&
A.isStringAttribute()) {
7535 AddingAttrs = RemovingAttrs =
true;
7538 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7539 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7541 RemovingAttrs =
true;
7542 if (
A.getValueAsString() ==
"true") {
7543 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7552 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7555 if (
A.getValueAsBool()) {
7556 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7562 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7563 RemovingAttrs =
true;
7570 bool HandleDenormalMode =
false;
7572 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7575 DenormalFPMath = ParsedMode;
7577 AddingAttrs = RemovingAttrs =
true;
7578 HandleDenormalMode =
true;
7582 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7586 DenormalFPMathF32 = ParsedMode;
7588 AddingAttrs = RemovingAttrs =
true;
7589 HandleDenormalMode =
true;
7593 if (HandleDenormalMode)
7594 AttrsToAdd.addDenormalFPEnvAttr(
7598 F.removeFnAttrs(AttrsToRemove);
7601 F.addFnAttrs(AttrsToAdd);
7607 if (!
F.hasFnAttribute(FnAttrName)) {
7608 F.addFnAttr(FnAttrName,
Value);
7610 <<
"\", function: " <<
F.getName() <<
"\n");
7618 if (!
F.hasFnAttribute(FnAttrName)) {
7620 F.addFnAttr(FnAttrName);
7622 <<
", function: " <<
F.getName() <<
"\n");
7625 auto A =
F.getFnAttribute(FnAttrName);
7626 if (
"false" ==
A.getValueAsString()) {
7627 F.removeFnAttr(FnAttrName);
7629 <<
"=\"false\", function: " <<
F.getName() <<
"\n");
7630 }
else if (
"true" ==
A.getValueAsString()) {
7631 F.removeFnAttr(FnAttrName);
7632 F.addFnAttr(FnAttrName);
7634 <<
"=\"true\", function: " <<
F.getName() <<
"\n");
7641 M.setModuleFlag(Behavior,
Key, Val);
7642 LLVM_DEBUG(
dbgs() <<
"Converted module flag: " <<
"{" << Behavior <<
", "
7643 <<
Key <<
", " << Val <<
"}\n");
7647 Triple T(M.getTargetTriple());
7648 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7651 uint64_t BTEValue = 0;
7652 uint64_t BPPLRValue = 0;
7653 uint64_t GCSValue = 0;
7654 uint64_t SRAValue = 0;
7655 uint64_t SRAALLValue = 0;
7656 uint64_t SRABKeyValue = 0;
7658 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7662 if (
Op->getNumOperands() != 3)
7671 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7672 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7673 : IDStr ==
"guarded-control-stack" ? &GCSValue
7674 : IDStr ==
"sign-return-address" ? &SRAValue
7675 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7676 : IDStr ==
"sign-return-address-with-bkey"
7682 *ValPtr = CI->getZExtValue();
7686 LLVM_DEBUG(
dbgs() <<
"Found module flag: " << IDStr <<
"(" << *ValPtr
7691 bool BTE = BTEValue == 1;
7692 bool BPPLR = BPPLRValue == 1;
7693 bool GCS = GCSValue == 1;
7694 bool SRA = SRAValue == 1;
7697 if (SRA && SRAALLValue == 1)
7698 SignTypeValue =
"all";
7701 if (SRA && SRABKeyValue == 1)
7702 SignKeyValue =
"b_key";
7704 for (
Function &
F : M.getFunctionList()) {
7705 if (
F.isDeclaration())
7712 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7713 A.isValid() &&
"none" ==
A.getValueAsString()) {
7714 F.removeFnAttr(
"sign-return-address");
7715 F.removeFnAttr(
"sign-return-address-key");
7731 if (SRAALLValue == 1)
7733 if (SRABKeyValue == 1) {
7762 if (
T->getNumOperands() < 1)
7767 if (S->getString().starts_with(
"llvm.vectorizer."))
7773 StringRef OldPrefix =
"llvm.vectorizer.";
7776 if (OldTag ==
"llvm.vectorizer.unroll")
7788 if (
T->getNumOperands() < 1)
7800 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7813 Ops.reserve(
T->getNumOperands());
7814 Ops.push_back(NewTag);
7815 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7816 Ops.push_back(
T->getOperand(
I));
7833 if (
T->isDistinct()) {
7834 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7846 Ops.reserve(
T->getNumOperands());
7857 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7858 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7859 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7862 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7864 auto I =
DL.find(
"-n64-");
7866 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7871 std::string Res =
DL.str();
7874 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7875 Res.append(Res.empty() ?
"G1" :
"-G1");
7883 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7884 Res.append(
"-ni:7:8:9");
7886 if (
DL.ends_with(
"ni:7"))
7888 if (
DL.ends_with(
"ni:7:8"))
7893 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7894 Res.append(
"-p7:160:256:256:32");
7895 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7896 Res.append(
"-p8:128:128:128:48");
7897 constexpr StringRef OldP8(
"-p8:128:128-");
7898 if (
DL.contains(OldP8))
7899 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7900 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7901 Res.append(
"-p9:192:256:256:32");
7906 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7907 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7908 Res.append((
"-" + AS +
":32:32").str());
7913 if (!
DL.contains(
"m:e"))
7914 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7919 if (
T.isSystemZ() && !
DL.empty()) {
7921 if (!
DL.contains(
"-S64"))
7922 return "E-S64" +
DL.drop_front(1).str();
7926 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7929 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7930 if (!
DL.contains(AddrSpaces)) {
7932 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7933 if (R.match(Res, &
Groups))
7939 if (
T.isAArch64()) {
7941 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7942 Res.append(
"-Fn32");
7943 AddPtr32Ptr64AddrSpaces();
7947 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7951 std::string I64 =
"-i64:64";
7952 std::string I128 =
"-i128:128";
7954 size_t Pos = Res.find(I64);
7955 if (Pos !=
size_t(-1))
7956 Res.insert(Pos + I64.size(), I128);
7960 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
7961 size_t Pos = Res.find(
"-S128");
7964 Res.insert(Pos,
"-f64:32:64");
7969 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
7971 size_t Pos = Res.
find(p3232);
7973 Res.insert(Pos + p3232.
size(),
"-Fi8");
7979 AddPtr32Ptr64AddrSpaces();
7987 if (!
T.isOSIAMCU()) {
7988 std::string I128 =
"-i128:128";
7991 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
7992 if (R.match(Res, &
Groups))
8000 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
8002 auto I =
Ref.find(
"-f80:32-");
8004 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
8012 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
8015 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
8016 B.removeAttribute(
"no-frame-pointer-elim");
8018 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
8020 if (FramePointer !=
"all")
8021 FramePointer =
"non-leaf";
8022 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
8024 if (!FramePointer.
empty())
8025 B.addAttribute(
"frame-pointer", FramePointer);
8027 A =
B.getAttribute(
"null-pointer-is-valid");
8030 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
8031 B.removeAttribute(
"null-pointer-is-valid");
8032 if (NullPointerIsValid)
8033 B.addAttribute(Attribute::NullPointerIsValid);
8036 A =
B.getAttribute(
"uniform-work-group-size");
8040 bool IsTrue = Val ==
"true";
8041 B.removeAttribute(
"uniform-work-group-size");
8043 B.addAttribute(
"uniform-work-group-size");
8054 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
static unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID, SmallVectorImpl< Type * > &OverloadTys)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI, StringRef Name, const Intrinsic::ID(&IIDs)[2][2])
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static constexpr Intrinsic::ID NVVMFAddIIDs[2][2]
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFPArithUpgrade(StringRef Name, const Intrinsic::ID(&IIDs)[2][2])
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static bool isOldLoopArgument(Metadata *MD)
static Value * upgradeARMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static void ConvertModuleFlag(Module &M, Module::ModFlagBehavior Behavior, StringRef Key, uint32_t Val)
static bool upgradeX86IntrinsicsWith8BitMask(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeVectorSplice(CallBase *CI, IRBuilder<> &Builder)
static Value * upgradeAMDGCNIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Value * upgradeMaskedLoad(IRBuilder<> &Builder, Value *Ptr, Value *Passthru, Value *Mask, bool Aligned)
static Metadata * unwrapMAVMetadataOp(CallBase *CI, unsigned Op)
Helper to unwrap Metadata MetadataAsValue operands, such as the Value field.
static bool upgradeX86BF16Intrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, StringRef Name, Function *&NewFn)
static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn, IRBuilder<> &Builder)
static Value * getX86MaskVec(IRBuilder<> &Builder, Value *Mask, unsigned NumElts)
static Value * emitX86ScalarSelect(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn)
static Value * upgradeX86ConcatShift(IRBuilder<> &Builder, CallBase &CI, bool IsShiftRight, bool ZeroMask)
static void rename(GlobalValue *GV)
static bool upgradePTESTIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradeX86BF16DPIntrinsic(Function *F, Intrinsic::ID IID, Function *&NewFn)
#define NVVM_TMA_G2S_MODES(M)
static Value * upgradeMaskedCompare(IRBuilder<> &Builder, CallBase &CI, unsigned CC, bool Signed)
static Value * upgradeX86BinaryIntrinsics(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static Value * upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Value * upgradeX86MaskedShift(IRBuilder<> &Builder, CallBase &CI, Intrinsic::ID IID)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static constexpr Intrinsic::ID NVVMFMulIIDs[2][2]
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static void reportFatalUsageErrorWithCI(StringRef reason, CallBase *CI)
static Value * upgradeMaskedStore(IRBuilder<> &Builder, Value *Ptr, Value *Data, Value *Mask, bool Aligned)
static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F, StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F, StringRef Name)
static Value * upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_CASE(ID_SUFFIX, NAME)
static bool upgradeX86MultiplyAddWords(Function *F, Intrinsic::ID IID, Function *&NewFn)
static bool upgradePtrauthInitFiniArrays(Module &M)
static Value * upgradeX86IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
static FCmpInst::Predicate getVPFPPredicateFromMD(const Value *Op)
static GCRegistry::Add< ShadowStackGC > C("shadow-stack", "Very portable GC for uncooperative code generators")
static GCRegistry::Add< ErlangGC > A("erlang", "erlang-compatible garbage collector")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
This file contains constants used for implementing Dwarf debug support.
Module.h This file contains the declarations for the Module class.
const AbstractManglingParser< Derived, Alloc >::OperatorInfo AbstractManglingParser< Derived, Alloc >::Ops[]
static bool isZero(Value *V, const DataLayout &DL, DominatorTree *DT, AssumptionCache *AC)
NVPTX address space definition.
This file contains the definitions of the enumerations and flags associated with NVVM Intrinsics,...
static bool contains(SmallPtrSetImpl< ConstantExpr * > &Cache, ConstantExpr *Expr, Constant *C)
This file implements the StringSwitch template, which mimics a switch() statement whose cases are str...
static SymbolRef::Type getType(const Symbol *Sym)
LocallyHashedType DenseMapInfo< LocallyHashedType >::Empty
static const X86InstrFMA3Group Groups[]
Class for arbitrary precision integers.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
Class to represent array types.
static LLVM_ABI ArrayType * get(Type *ElementType, uint64_t NumElements)
This static method is the primary way to construct an ArrayType.
Type * getElementType() const
an instruction that atomically reads a memory location, combines it with another value,...
void setVolatile(bool V)
Specify whether this is a volatile RMW or not.
BinOp
This enumeration lists the possible modifications atomicrmw can make.
@ USubCond
Subtract only if no unsigned overflow.
@ Min
*p = old <signed v ? old : v
@ USubSat
*p = usub.sat(old, v) usub.sat matches the behavior of llvm.usub.sat.
@ UIncWrap
Increment one up to a maximum value.
@ Max
*p = old >signed v ? old : v
@ FMin
*p = minnum(old, v) minnum matches the behavior of llvm.minnum.
@ FMax
*p = maxnum(old, v) maxnum matches the behavior of llvm.maxnum.
@ UDecWrap
Decrement one until a minimum value or zero.
bool isFloatingPointOperation() const
This class stores enough information to efficiently remove some attributes from an existing AttrBuild...
AttributeMask & addAttribute(Attribute::AttrKind Val)
Add an attribute to the mask.
Functions, function parameters, and return types can have attributes to indicate how they should be t...
static LLVM_ABI Attribute getWithStackAlignment(LLVMContext &Context, Align Alignment)
static LLVM_ABI Attribute get(LLVMContext &Context, AttrKind Kind, uint64_t Val=0)
Return a uniquified Attribute object.
Base class for all callable instructions (InvokeInst and CallInst) Holds everything related to callin...
void setCallingConv(CallingConv::ID CC)
LLVM_ABI void getOperandBundlesAsDefs(SmallVectorImpl< OperandBundleDef > &Defs) const
Return the list of operand bundles attached to this instruction as a vector of OperandBundleDefs.
Function * getCalledFunction() const
Returns the function called, or null if this is an indirect function invocation or the function signa...
CallingConv::ID getCallingConv() const
Value * getCalledOperand() const
void setAttributes(AttributeList A)
Set the attributes for this call.
Value * getArgOperand(unsigned i) const
FunctionType * getFunctionType() const
LLVM_ABI Intrinsic::ID getIntrinsicID() const
Returns the intrinsic ID of the intrinsic called or Intrinsic::not_intrinsic if the called function i...
iterator_range< User::op_iterator > args()
Iteration adapter for range-for loops.
void setCalledOperand(Value *V)
unsigned arg_size() const
AttributeList getAttributes() const
Return the attributes for this call.
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
This class represents a function call, abstracting a target machine's calling convention.
void setTailCallKind(TailCallKind TCK)
static LLVM_ABI CastInst * Create(Instruction::CastOps, Value *S, Type *Ty, const Twine &Name="", InsertPosition InsertBefore=nullptr)
Provides a way to construct any of the CastInst subclasses using an opcode instead of the subclass's ...
static LLVM_ABI bool castIsValid(Instruction::CastOps op, Type *SrcTy, Type *DstTy)
This method can be used to determine if a cast from SrcTy to DstTy using Opcode op is valid or not.
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI ConstantAggregateZero * get(Type *Ty)
static LLVM_ABI Constant * get(ArrayType *T, ArrayRef< Constant * > V)
static LLVM_ABI Constant * getIntToPtr(Constant *C, Type *Ty, bool OnlyIfReduced=false)
static LLVM_ABI Constant * getPointerCast(Constant *C, Type *Ty)
Create a BitCast, AddrSpaceCast, or a PtrToInt cast constant expression.
static LLVM_ABI Constant * getPtrToInt(Constant *C, Type *Ty, bool OnlyIfReduced=false)
This is the shared class of boolean and integer constants.
bool isZero() const
This is just a convenience method to make client code smaller for a common code.
uint64_t getZExtValue() const
Return the constant as a 64-bit unsigned integer value after it has been zero extended as appropriate...
static LLVM_ABI ConstantPointerNull * get(PointerType *T)
Static factory methods - Return objects of the specified value.
static LLVM_ABI Constant * get(StructType *T, ArrayRef< Constant * > V)
StructType * getType() const
Specialization - reduce amount of casting.
static LLVM_ABI ConstantTokenNone * get(LLVMContext &Context)
Return the ConstantTokenNone.
This is an important base class in LLVM.
static LLVM_ABI Constant * getAllOnesValue(Type *Ty)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
static LLVM_ABI DIExpression * append(const DIExpression *Expr, ArrayRef< uint64_t > Ops)
Append the opcodes Ops to DIExpr.
A parsed version of the target data layout string in and methods for querying it.
static LLVM_ABI DbgLabelRecord * createUnresolvedDbgLabelRecord(MDNode *Label)
For use during parsing; creates a DbgLabelRecord from as-of-yet unresolved MDNodes.
Base class for non-instruction debug metadata records that have positions within IR.
void setDebugLoc(DebugLoc Loc)
static LLVM_ABI DbgVariableRecord * createUnresolvedDbgVariableRecord(LocationType Type, Metadata *Val, MDNode *Variable, MDNode *Expression, MDNode *AssignID, Metadata *Address, MDNode *AddressExpression)
Used to create DbgVariableRecords during parsing, where some metadata references may still be unresol...
Convenience struct for specifying and reasoning about fast-math flags.
void setApproxFunc(bool B=true)
static LLVM_ABI FixedVectorType * get(Type *ElementType, unsigned NumElts)
Class to represent function types.
unsigned getNumParams() const
Return the number of fixed parameters this function type requires.
Type * getParamType(unsigned i) const
Parameter type accessors.
Type * getReturnType() const
static LLVM_ABI FunctionType * get(Type *Result, ArrayRef< Type * > Params, bool isVarArg)
This static method is the primary way of constructing a FunctionType.
static Function * Create(FunctionType *Ty, LinkageTypes Linkage, unsigned AddrSpace, const Twine &N="", Module *M=nullptr)
FunctionType * getFunctionType() const
Returns the FunctionType for me.
Intrinsic::ID getIntrinsicID() const LLVM_READONLY
getIntrinsicID - This method returns the ID number of the specified function, or Intrinsic::not_intri...
const Function & getFunction() const
void eraseFromParent()
eraseFromParent - This method unlinks 'this' from the containing module and deletes it.
Type * getReturnType() const
Returns the type of the ret val.
Argument * getArg(unsigned i) const
static LLVM_ABI GUID getGUIDAssumingExternalLinkage(StringRef GlobalName)
Return a 64-bit global unique ID constructed from the name of a global symbol.
LinkageTypes getLinkage() const
uint64_t GUID
Declare a type to represent a global unique identifier for a global value.
static StringRef dropLLVMManglingEscape(StringRef Name)
If the given string begins with the GlobalValue name mangling escape character '\1',...
Module * getParent()
Get the module that this global value is contained inside of...
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
op_iterator op_end() const
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
unsigned getNumOperands() const
Return number of MDNode operands.
op_iterator op_begin() const
LLVMContext & getContext() const
Tracking metadata reference owned by Metadata.
LLVM_ABI StringRef getString() const
static LLVM_ABI MDString * get(LLVMContext &Context, StringRef Str)
static MDTuple * get(LLVMContext &Context, ArrayRef< Metadata * > MDs)
A Module instance is used to store all the information related to an LLVM module.
ModFlagBehavior
This enumeration defines the supported behaviors of module flags.
@ Override
Uses the specified value, regardless of the behavior or value of the other module.
@ Error
Emits an error if two values disagree, otherwise the resulting value is that of the operands.
@ Min
Takes the min of the two values, which are required to be integers.
@ Max
Takes the max of the two values, which are required to be integers.
LLVM_ABI void setOperand(unsigned I, MDNode *New)
LLVM_ABI MDNode * getOperand(unsigned i) const
LLVM_ABI unsigned getNumOperands() const
LLVM_ABI void clearOperands()
Drop all references to this node's operands.
iterator_range< op_iterator > operands()
LLVM_ABI void addOperand(MDNode *M)
ArrayRef< InputTy > inputs() const
static LLVM_ABI PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
LLVM_ABI bool match(StringRef String, SmallVectorImpl< StringRef > *Matches=nullptr, std::string *Error=nullptr) const
matches - Match the regex against a given String.
static LLVM_ABI ScalableVectorType * get(Type *ElementType, unsigned MinNumElts)
ArrayRef< int > getShuffleMask() const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
SmallString - A SmallString is just a SmallVector with methods and accessors that make it work better...
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
reference emplace_back(ArgTypes &&... Args)
void append(ItTy in_start, ItTy in_end)
Add the specified range to the end of the SmallVector.
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
An instruction for storing to memory.
A wrapper around a string literal that serves as a proxy for constructing global tables of StringRefs...
Represent a constant reference to a string, i.e.
std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
static constexpr size_t npos
constexpr StringRef substr(size_t Start, size_t N=npos) const
Return a reference to the substring from [Start, Start + N).
bool starts_with(StringRef Prefix) const
Check if this string starts with the given Prefix.
constexpr bool empty() const
Check if the string is empty.
StringRef drop_front(size_t N=1) const
Return a StringRef equal to 'this' but with the first N elements dropped.
constexpr size_t size() const
Get the string size.
size_t find(char C, size_t From=0) const
Search for the first character C in the string.
StringRef trim(char Char) const
Return string with consecutive Char characters starting from the left and right removed.
bool consume_front(char Prefix)
Returns true if this StringRef has the given prefix and removes that prefix.
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 setMetadata(unsigned KindID, MDNode *Node)
Set a particular kind of metadata attachment.
LLVM_ABI void setName(const Twine &Name)
Change the name of the value.
LLVM_ABI void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
LLVMContext & getContext() const
All values hold a context through their type.
iterator_range< user_iterator > users()
LLVM_ABI const Value * stripPointerCasts() const
Strip off pointer casts, all-zero GEPs and address space casts.
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
LLVM_ABI void takeName(Value *V)
Transfer the name from V to this value.
Base class of all SIMD vector types.
static VectorType * getInteger(VectorType *VTy)
This static method gets a VectorType with the same number of elements as the input type,...
static LLVM_ABI VectorType * get(Type *ElementType, ElementCount EC)
This static method is the primary way to construct an VectorType.
constexpr ScalarTy getFixedValue() const
const ParentTy * getParent() const
self_iterator getIterator()
A raw_ostream that writes to an SmallVector or SmallString.
StringRef str() const
Return a StringRef for the vector contents.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
@ LOCAL_ADDRESS
Address space for local memory.
@ FLAT_ADDRESS
Address space for flat memory.
@ PRIVATE_ADDRESS
Address space for private memory.
@ PTX_Kernel
Call to a PTX kernel. Passes all arguments in parameter space.
std::optional< ABIType > parseABIType(StringRef S)
Parse the string spelling used by the "float-abi" IR module flag into an ABIType.
LLVM_ABI std::optional< Function * > remangleIntrinsicFunction(Function *F)
LLVM_ABI Function * getOrInsertDeclaration(Module *M, ID id, ArrayRef< Type * > OverloadTys={})
Look up the Function declaration of the intrinsic id in the Module M.
LLVM_ABI AttributeList getAttributes(LLVMContext &C, ID id, FunctionType *FT)
Return the attributes for an intrinsic.
LLVM_ABI bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
LLVM_ABI MDNode * UpgradeTBAAStructNode(MDNode &TBAAStructNode)
If the given !tbaa.struct node has old-style scalar field tags, return an equivalent node with each f...
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_ATTRIBUTE_VISIBILITY_DEFAULT AnalysisKey InnerAnalysisManagerProxy< AnalysisManagerT, IRUnitT, ExtraArgTs... >::Key
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
const IROptions & getIROptions(const LLVMContext &C)
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
bool is_contained(R &&Range, const E &Element)
Returns true if Element is found in Range.
LLVM_ABI bool UpgradeDebugInfo(Module &M)
Check the debug info version number, if it is out-dated, drop the debug info.
LLVM_ABI void UpgradeFunctionAttributes(Function &F)
Correct any IR that is relying on old function attribute behavior.
LLVM_ABI MDNode * UpgradeTBAANode(MDNode &TBAANode)
If the given TBAA tag uses the scalar TBAA format, create a new node corresponding to the upgrade to ...
LLVM_ABI void UpgradeARCRuntime(Module &M)
Convert calls to ARC runtime functions to intrinsic calls and upgrade the old retain release marker t...
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.