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)
1426 .
Case(
"tcgen05.alloc.shared.cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1427 .
Case(
"tcgen05.alloc.shared.cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1432 if (Name.consume_front(
"fmax."))
1434 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1435 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1436 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1437 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1438 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1439 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1440 .
Case(
"ftz.nan.xorsign.abs.bf16",
1441 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1442 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1443 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1444 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1445 .
Case(
"ftz.xorsign.abs.bf16x2",
1446 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1447 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1448 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1449 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1450 .
Case(
"nan.xorsign.abs.bf16x2",
1451 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1452 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1453 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1456 if (Name.consume_front(
"fmin."))
1458 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1459 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1460 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1461 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1462 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1463 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1464 .
Case(
"ftz.nan.xorsign.abs.bf16",
1465 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1466 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1467 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1468 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1469 .
Case(
"ftz.xorsign.abs.bf16x2",
1470 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1471 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1472 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1473 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1474 .
Case(
"nan.xorsign.abs.bf16x2",
1475 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1476 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1477 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1480 if (Name.consume_front(
"neg."))
1482 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1483 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1492 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1497 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1500 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1503 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1504 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1514 {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat,
1518 {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat,
1521 {Intrinsic::nvvm_fma, Intrinsic::nvvm_fma_sat, Intrinsic::nvvm_fma_relu},
1522 {Intrinsic::nvvm_fma_ftz, Intrinsic::nvvm_fma_ftz_sat,
1523 Intrinsic::nvvm_fma_ftz_relu}};
1529 bool IsFTZ,
bool IsSat,
1534 .
Case(
"f", !IsRelu && !(IsSat && IsFMul))
1535 .
Case(
"d", !IsFTZ && !IsSat && !IsRelu)
1536 .
Case(
"f16", IsRN && (IsFMA || IsSat))
1537 .
Case(
"v2f16", IsRN && !IsFMA && IsSat)
1538 .
Case(
"f16x2", IsRN && IsFMA)
1539 .
Cases({
"bf16",
"bf16x2"}, IsRN && IsFMA && !IsFTZ && !IsSat)
1543static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1545 auto [Modifiers,
Type] = Name.rsplit(
'.');
1555 return std::nullopt;
1557 StringRef Rest = Modifiers.drop_front(2);
1562 return std::nullopt;
1564 const Intrinsic::ID IID = IIDs[IsFTZ][IsSat ? 1 : IsRelu ? 2 : 0];
1566 return std::nullopt;
1572 return std::nullopt;
1578 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1581 return Intrinsic::nvvm_mbarrier_init;
1585 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1586 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1587 Name.consume_front(
"param");
1591 if (!Name.consume_front(
"vp."))
1620 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1621 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1628 if (!Name.consume_front(
"vp."))
1648 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1649 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1654 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1666 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1677 if (Name.starts_with(
"to.fp16")) {
1681 FuncTy->getReturnType());
1684 if (Name.starts_with(
"from.fp16")) {
1688 FuncTy->getReturnType());
1698 if (Defaults.empty())
1701 unsigned FullArgCount = FirstDefault + Defaults.size();
1704 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1707 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1709 NumMissingTrailingParams))
1712 return FullArgCount;
1719 unsigned FullArgCount =
1721 if (FullArgCount == 0)
1727 "total number of default args does not match intrinsic signature");
1732 bool CanUpgradeDebugIntrinsicsToRecords) {
1733 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1738 if (!Name.consume_front(
"llvm.") || Name.empty())
1744 bool IsArm = Name.consume_front(
"arm.");
1745 if (IsArm || Name.consume_front(
"aarch64.")) {
1751 if (Name.consume_front(
"amdgcn.")) {
1752 if (Name ==
"alignbit") {
1755 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1759 if (Name.consume_front(
"atomic.")) {
1760 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1761 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1770 if (Name.starts_with(
"addrspacecast.nonnull")) {
1777 switch (
F->getIntrinsicID()) {
1781 case Intrinsic::amdgcn_asyncmark:
1782 case Intrinsic::amdgcn_wait_asyncmark: {
1784 if (
F->arg_size() == (ID == Intrinsic::amdgcn_asyncmark ? 0u : 1u)) {
1792 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1793 if (
F->arg_size() == 7) {
1798 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1799 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1800 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1801 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1802 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1803 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1804 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1805 if (
F->arg_size() == 8) {
1812 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1813 Name.consume_front(
"flat.atomic.")) {
1814 if (Name.starts_with(
"fadd") ||
1816 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1817 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1825 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1830 if (Name.starts_with(
"ldexp.")) {
1833 F->getParent(), Intrinsic::ldexp,
1834 {F->getReturnType(), F->getArg(1)->getType()});
1843 if (
F->arg_size() == 1) {
1844 if (Name.consume_front(
"convert.")) {
1858 F->arg_begin()->getType());
1864 if (Name ==
"coro.end" &&
1865 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1866 CoroEndID = Intrinsic::coro_end;
1867 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1868 CoroEndID = Intrinsic::coro_end_async;
1879 if (Name.consume_front(
"dbg.")) {
1881 if (CanUpgradeDebugIntrinsicsToRecords) {
1882 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1883 Name ==
"declare" || Name ==
"label") {
1892 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1895 Intrinsic::dbg_value);
1902 if (Name.consume_front(
"experimental.vector.")) {
1908 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1909 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1910 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1911 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1912 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1914 Intrinsic::vector_partial_reduce_add)
1917 const auto *FT =
F->getFunctionType();
1919 if (ID == Intrinsic::vector_extract ||
1920 ID == Intrinsic::vector_interleave2)
1923 if (ID != Intrinsic::vector_interleave2)
1925 if (ID == Intrinsic::vector_insert ||
1926 ID == Intrinsic::vector_partial_reduce_add)
1934 if (Name.consume_front(
"reduce.")) {
1936 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1937 if (R.match(Name, &
Groups))
1939 .
Case(
"add", Intrinsic::vector_reduce_add)
1940 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1941 .
Case(
"and", Intrinsic::vector_reduce_and)
1942 .
Case(
"or", Intrinsic::vector_reduce_or)
1943 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1944 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1945 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1946 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1947 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1948 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1949 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1954 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1959 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1960 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1965 auto Args =
F->getFunctionType()->params();
1967 {Args[V2 ? 1 : 0]});
1973 if (Name.consume_front(
"splice"))
1977 if (Name.consume_front(
"experimental.stepvector.")) {
1981 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1986 if (Name.starts_with(
"flt.rounds")) {
1989 Intrinsic::get_rounding);
1994 if (Name.starts_with(
"invariant.group.barrier")) {
1996 auto Args =
F->getFunctionType()->params();
1997 Type* ObjectPtr[1] = {Args[0]};
2000 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
2005 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
2006 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
2007 if (IsLifetimeStart || IsLifetimeEnd) {
2008 if (
F->arg_size() == 2) {
2009 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
2010 : Intrinsic::lifetime_end;
2015 F->getArg(1)->getType());
2017 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2037 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2038 .StartsWith(
"memmove.", Intrinsic::memmove)
2040 if (
F->arg_size() == 5) {
2044 F->getFunctionType()->params().slice(0, 3);
2050 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2053 const auto *FT =
F->getFunctionType();
2054 Type *ParamTypes[2] = {
2055 FT->getParamType(0),
2059 Intrinsic::memset, ParamTypes);
2065 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2066 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2067 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2068 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2070 if (MaskedID &&
F->arg_size() == 4) {
2072 if (MaskedID == Intrinsic::masked_load ||
2073 MaskedID == Intrinsic::masked_gather) {
2075 F->getParent(), MaskedID,
2076 {F->getReturnType(), F->getArg(0)->getType()});
2080 F->getParent(), MaskedID,
2081 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2087 if (Name.consume_front(
"nvvm.")) {
2089 if (
F->arg_size() == 1) {
2092 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2093 .Case(
"clz.i", Intrinsic::ctlz)
2094 .
Case(
"popc.i", Intrinsic::ctpop)
2098 {F->getReturnType()});
2101 }
else if (
F->arg_size() == 2) {
2104 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2105 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2106 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2107 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2108 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2109 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2113 {F->getReturnType()});
2150 F->getParent(), IID,
F->getReturnType(),
2151 F->getFunctionType()->params());
2160 {F->getArg(0)->getType()});
2206 F->getArg(0)->getType());
2214 bool Expand =
false;
2215 if (Name.consume_front(
"abs."))
2218 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2219 else if (Name.consume_front(
"fabs."))
2221 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2222 else if (Name.consume_front(
"add."))
2225 else if (Name.consume_front(
"mul."))
2228 else if (Name.consume_front(
"ex2.approx."))
2231 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2232 else if (Name.consume_front(
"fma."))
2235 else if (Name.consume_front(
"atomic.load."))
2244 else if (Name.consume_front(
"atomic."))
2259 else if (Name.consume_front(
"bitcast."))
2262 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2263 else if (Name.consume_front(
"rotate."))
2265 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2266 else if (Name.consume_front(
"ptr.gen.to."))
2269 else if (Name.consume_front(
"ptr."))
2272 else if (Name.consume_front(
"ldg.global."))
2274 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2275 Name.starts_with(
"p."));
2278 .
Case(
"barrier0",
true)
2279 .
Case(
"barrier.n",
true)
2280 .
Case(
"barrier.sync.cnt",
true)
2281 .
Case(
"barrier.sync",
true)
2282 .
Case(
"barrier",
true)
2283 .
Case(
"bar.sync",
true)
2284 .
Case(
"barrier0.popc",
true)
2285 .
Case(
"barrier0.and",
true)
2286 .
Case(
"barrier0.or",
true)
2287 .
Case(
"clz.ll",
true)
2288 .
Case(
"popc.ll",
true)
2290 .
Case(
"swap.lo.hi.b64",
true)
2291 .
Case(
"tanh.approx.f32",
true)
2303 if (Name.starts_with(
"objectsize.")) {
2304 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2305 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2308 Intrinsic::objectsize, Tys);
2315 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2318 F->getParent(), Intrinsic::ptr_annotation,
2319 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2325 if (Name.consume_front(
"riscv.")) {
2328 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2329 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2330 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2331 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2334 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2347 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2348 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2357 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2358 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2359 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2360 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2365 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2374 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2376 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2385 if (Name ==
"stackprotectorcheck") {
2389 if (Name.starts_with(
"strip.invariant.group")) {
2394 F->getParent(), Intrinsic::launder_invariant_group,
2395 F->getReturnType());
2401 if (Name ==
"thread.pointer") {
2403 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2409 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2412 F->getParent(), Intrinsic::var_annotation,
2413 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2416 if (Name.consume_front(
"vector.splice")) {
2417 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2427 if (Name.consume_front(
"wasm.")) {
2430 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2431 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2432 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2437 F->getReturnType());
2441 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2443 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2445 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2464 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2474 std::string
Name =
F->getName().str();
2477 Name,
F->getParent());
2488 if (Result != std::nullopt) {
2505 bool CanUpgradeDebugIntrinsicsToRecords) {
2525 GV->
getName() ==
"llvm.global_dtors")) ||
2539 unsigned N =
Init->getNumOperands();
2540 std::vector<Constant *> NewCtors(
N);
2541 for (
unsigned i = 0; i !=
N; ++i) {
2544 Ctor->getAggregateElement(1),
2558 unsigned NumElts = ResultTy->getNumElements() * 8;
2562 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2572 for (
unsigned l = 0; l != NumElts; l += 16)
2573 for (
unsigned i = 0; i != 16; ++i) {
2574 unsigned Idx = NumElts + i - Shift;
2576 Idx -= NumElts - 16;
2577 Idxs[l + i] = Idx + l;
2580 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2584 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2592 unsigned NumElts = ResultTy->getNumElements() * 8;
2596 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2606 for (
unsigned l = 0; l != NumElts; l += 16)
2607 for (
unsigned i = 0; i != 16; ++i) {
2608 unsigned Idx = i + Shift;
2610 Idx += NumElts - 16;
2611 Idxs[l + i] = Idx + l;
2614 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2618 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2626 Mask = Builder.CreateBitCast(Mask, MaskTy);
2632 for (
unsigned i = 0; i != NumElts; ++i)
2634 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2645 if (
C->isAllOnesValue())
2650 return Builder.CreateSelect(Mask, Op0, Op1);
2657 if (
C->isAllOnesValue())
2661 Mask->getType()->getIntegerBitWidth());
2662 Mask = Builder.CreateBitCast(Mask, MaskTy);
2663 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2664 return Builder.CreateSelect(Mask, Op0, Op1);
2677 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2678 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2683 ShiftVal &= (NumElts - 1);
2692 if (ShiftVal > 16) {
2700 for (
unsigned l = 0; l < NumElts; l += 16) {
2701 for (
unsigned i = 0; i != 16; ++i) {
2702 unsigned Idx = ShiftVal + i;
2703 if (!IsVALIGN && Idx >= 16)
2704 Idx += NumElts - 16;
2705 Indices[l + i] = Idx + l;
2710 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2716 bool ZeroMask,
bool IndexForm) {
2719 unsigned EltWidth = Ty->getScalarSizeInBits();
2720 bool IsFloat = Ty->isFPOrFPVectorTy();
2722 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2723 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2724 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2725 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2726 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2727 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2728 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2729 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2730 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2731 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2732 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2733 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2734 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2735 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2736 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2737 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2738 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2739 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2740 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2741 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2742 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2743 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2744 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2745 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2746 else if (VecWidth == 128 && EltWidth == 16)
2747 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2748 else if (VecWidth == 256 && EltWidth == 16)
2749 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2750 else if (VecWidth == 512 && EltWidth == 16)
2751 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2752 else if (VecWidth == 128 && EltWidth == 8)
2753 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2754 else if (VecWidth == 256 && EltWidth == 8)
2755 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2756 else if (VecWidth == 512 && EltWidth == 8)
2757 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2768 Value *
V = Builder.CreateIntrinsic(IID, Args);
2780 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2791 bool IsRotateRight) {
2801 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2802 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2805 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2806 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2851 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2856 bool IsShiftRight,
bool ZeroMask) {
2870 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2871 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2874 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2875 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2890 const Align Alignment =
2892 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2897 if (
C->isAllOnesValue())
2898 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2903 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2909 const Align Alignment =
2918 if (
C->isAllOnesValue())
2919 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2924 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2930 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2931 {Op0, Builder.getInt1(
false)});
2946 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2947 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2948 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2949 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2950 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2953 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2954 LHS = Builder.CreateAnd(
LHS, Mask);
2955 RHS = Builder.CreateAnd(
RHS, Mask);
2972 if (!
C || !
C->isAllOnesValue())
2973 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2978 for (
unsigned i = 0; i != NumElts; ++i)
2980 for (
unsigned i = NumElts; i != 8; ++i)
2981 Indices[i] = NumElts + i % NumElts;
2982 Vec = Builder.CreateShuffleVector(Vec,
2986 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
2998 }
else if (
CC == 7) {
3034 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3035 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3037 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3038 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3047 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3053 Name = Name.substr(12);
3058 if (Name.starts_with(
"max.p")) {
3059 if (VecWidth == 128 && EltWidth == 32)
3060 IID = Intrinsic::x86_sse_max_ps;
3061 else if (VecWidth == 128 && EltWidth == 64)
3062 IID = Intrinsic::x86_sse2_max_pd;
3063 else if (VecWidth == 256 && EltWidth == 32)
3064 IID = Intrinsic::x86_avx_max_ps_256;
3065 else if (VecWidth == 256 && EltWidth == 64)
3066 IID = Intrinsic::x86_avx_max_pd_256;
3069 }
else if (Name.starts_with(
"min.p")) {
3070 if (VecWidth == 128 && EltWidth == 32)
3071 IID = Intrinsic::x86_sse_min_ps;
3072 else if (VecWidth == 128 && EltWidth == 64)
3073 IID = Intrinsic::x86_sse2_min_pd;
3074 else if (VecWidth == 256 && EltWidth == 32)
3075 IID = Intrinsic::x86_avx_min_ps_256;
3076 else if (VecWidth == 256 && EltWidth == 64)
3077 IID = Intrinsic::x86_avx_min_pd_256;
3080 }
else if (Name.starts_with(
"pshuf.b.")) {
3081 if (VecWidth == 128)
3082 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3083 else if (VecWidth == 256)
3084 IID = Intrinsic::x86_avx2_pshuf_b;
3085 else if (VecWidth == 512)
3086 IID = Intrinsic::x86_avx512_pshuf_b_512;
3089 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3090 if (VecWidth == 128)
3091 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3092 else if (VecWidth == 256)
3093 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3094 else if (VecWidth == 512)
3095 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3098 }
else if (Name.starts_with(
"pmulh.w")) {
3099 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3100 "Unexpected intrinsic");
3103 }
else if (Name.starts_with(
"pmulhu.w")) {
3104 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3105 "Unexpected intrinsic");
3108 }
else if (Name.starts_with(
"pmaddw.d.")) {
3109 if (VecWidth == 128)
3110 IID = Intrinsic::x86_sse2_pmadd_wd;
3111 else if (VecWidth == 256)
3112 IID = Intrinsic::x86_avx2_pmadd_wd;
3113 else if (VecWidth == 512)
3114 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3117 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3118 if (VecWidth == 128)
3119 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3120 else if (VecWidth == 256)
3121 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3122 else if (VecWidth == 512)
3123 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3126 }
else if (Name.starts_with(
"packsswb.")) {
3127 if (VecWidth == 128)
3128 IID = Intrinsic::x86_sse2_packsswb_128;
3129 else if (VecWidth == 256)
3130 IID = Intrinsic::x86_avx2_packsswb;
3131 else if (VecWidth == 512)
3132 IID = Intrinsic::x86_avx512_packsswb_512;
3135 }
else if (Name.starts_with(
"packssdw.")) {
3136 if (VecWidth == 128)
3137 IID = Intrinsic::x86_sse2_packssdw_128;
3138 else if (VecWidth == 256)
3139 IID = Intrinsic::x86_avx2_packssdw;
3140 else if (VecWidth == 512)
3141 IID = Intrinsic::x86_avx512_packssdw_512;
3144 }
else if (Name.starts_with(
"packuswb.")) {
3145 if (VecWidth == 128)
3146 IID = Intrinsic::x86_sse2_packuswb_128;
3147 else if (VecWidth == 256)
3148 IID = Intrinsic::x86_avx2_packuswb;
3149 else if (VecWidth == 512)
3150 IID = Intrinsic::x86_avx512_packuswb_512;
3153 }
else if (Name.starts_with(
"packusdw.")) {
3154 if (VecWidth == 128)
3155 IID = Intrinsic::x86_sse41_packusdw;
3156 else if (VecWidth == 256)
3157 IID = Intrinsic::x86_avx2_packusdw;
3158 else if (VecWidth == 512)
3159 IID = Intrinsic::x86_avx512_packusdw_512;
3162 }
else if (Name.starts_with(
"vpermilvar.")) {
3163 if (VecWidth == 128 && EltWidth == 32)
3164 IID = Intrinsic::x86_avx_vpermilvar_ps;
3165 else if (VecWidth == 128 && EltWidth == 64)
3166 IID = Intrinsic::x86_avx_vpermilvar_pd;
3167 else if (VecWidth == 256 && EltWidth == 32)
3168 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3169 else if (VecWidth == 256 && EltWidth == 64)
3170 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3171 else if (VecWidth == 512 && EltWidth == 32)
3172 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3173 else if (VecWidth == 512 && EltWidth == 64)
3174 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3177 }
else if (Name ==
"cvtpd2dq.256") {
3178 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3179 }
else if (Name ==
"cvtpd2ps.256") {
3180 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3181 }
else if (Name ==
"cvttpd2dq.256") {
3182 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3183 }
else if (Name ==
"cvttps2dq.128") {
3184 IID = Intrinsic::x86_sse2_cvttps2dq;
3185 }
else if (Name ==
"cvttps2dq.256") {
3186 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3187 }
else if (Name.starts_with(
"permvar.")) {
3189 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3190 IID = Intrinsic::x86_avx2_permps;
3191 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3192 IID = Intrinsic::x86_avx2_permd;
3193 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3194 IID = Intrinsic::x86_avx512_permvar_df_256;
3195 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3196 IID = Intrinsic::x86_avx512_permvar_di_256;
3197 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3198 IID = Intrinsic::x86_avx512_permvar_sf_512;
3199 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3200 IID = Intrinsic::x86_avx512_permvar_si_512;
3201 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3202 IID = Intrinsic::x86_avx512_permvar_df_512;
3203 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3204 IID = Intrinsic::x86_avx512_permvar_di_512;
3205 else if (VecWidth == 128 && EltWidth == 16)
3206 IID = Intrinsic::x86_avx512_permvar_hi_128;
3207 else if (VecWidth == 256 && EltWidth == 16)
3208 IID = Intrinsic::x86_avx512_permvar_hi_256;
3209 else if (VecWidth == 512 && EltWidth == 16)
3210 IID = Intrinsic::x86_avx512_permvar_hi_512;
3211 else if (VecWidth == 128 && EltWidth == 8)
3212 IID = Intrinsic::x86_avx512_permvar_qi_128;
3213 else if (VecWidth == 256 && EltWidth == 8)
3214 IID = Intrinsic::x86_avx512_permvar_qi_256;
3215 else if (VecWidth == 512 && EltWidth == 8)
3216 IID = Intrinsic::x86_avx512_permvar_qi_512;
3219 }
else if (Name.starts_with(
"dbpsadbw.")) {
3220 if (VecWidth == 128)
3221 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3222 else if (VecWidth == 256)
3223 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3224 else if (VecWidth == 512)
3225 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3228 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3229 if (VecWidth == 128)
3230 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3231 else if (VecWidth == 256)
3232 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3233 else if (VecWidth == 512)
3234 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3237 }
else if (Name.starts_with(
"conflict.")) {
3238 if (Name[9] ==
'd' && VecWidth == 128)
3239 IID = Intrinsic::x86_avx512_conflict_d_128;
3240 else if (Name[9] ==
'd' && VecWidth == 256)
3241 IID = Intrinsic::x86_avx512_conflict_d_256;
3242 else if (Name[9] ==
'd' && VecWidth == 512)
3243 IID = Intrinsic::x86_avx512_conflict_d_512;
3244 else if (Name[9] ==
'q' && VecWidth == 128)
3245 IID = Intrinsic::x86_avx512_conflict_q_128;
3246 else if (Name[9] ==
'q' && VecWidth == 256)
3247 IID = Intrinsic::x86_avx512_conflict_q_256;
3248 else if (Name[9] ==
'q' && VecWidth == 512)
3249 IID = Intrinsic::x86_avx512_conflict_q_512;
3252 }
else if (Name.starts_with(
"pavg.")) {
3253 if (Name[5] ==
'b' && VecWidth == 128)
3254 IID = Intrinsic::x86_sse2_pavg_b;
3255 else if (Name[5] ==
'b' && VecWidth == 256)
3256 IID = Intrinsic::x86_avx2_pavg_b;
3257 else if (Name[5] ==
'b' && VecWidth == 512)
3258 IID = Intrinsic::x86_avx512_pavg_b_512;
3259 else if (Name[5] ==
'w' && VecWidth == 128)
3260 IID = Intrinsic::x86_sse2_pavg_w;
3261 else if (Name[5] ==
'w' && VecWidth == 256)
3262 IID = Intrinsic::x86_avx2_pavg_w;
3263 else if (Name[5] ==
'w' && VecWidth == 512)
3264 IID = Intrinsic::x86_avx512_pavg_w_512;
3273 Rep = Builder.CreateIntrinsic(IID, Args);
3284 if (AsmStr->find(
"mov\tfp") == 0 &&
3285 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3286 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3287 AsmStr->replace(Pos, 1,
";");
3295 assert(Result &&
"unsupported nvvm.{add,mul,fma}.* intrinsic");
3303 if (IsLegacyBF16Storage)
3304 Ty = Name.ends_with(
"bf16x2")
3306 :
static_cast<Type *
>(Builder.getBFloatTy());
3310 Args.push_back(IsLegacyBF16Storage ? Builder.CreateBitCast(Arg, Ty) : Arg);
3311 Args.push_back(Builder.getInt32(
static_cast<int>(
RoundingMode)));
3313 Value *Rep = Builder.CreateIntrinsic(Ty, IID, Args);
3314 return IsLegacyBF16Storage ? Builder.CreateBitCast(Rep, CI->
getType()) : Rep;
3319 Value *Rep =
nullptr;
3321 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3323 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3324 {Arg, Builder.getTrue()},
3326 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3327 Type *Ty = (Name ==
"abs.bf16")
3331 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3332 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3333 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3334 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3335 : Intrinsic::nvvm_fabs;
3336 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3337 }
else if (Name.consume_front(
"add.")) {
3340 }
else if (Name.consume_front(
"mul.")) {
3343 }
else if (Name.consume_front(
"ex2.approx.")) {
3345 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3346 : Intrinsic::nvvm_ex2_approx;
3347 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3348 }
else if (Name.consume_front(
"fma.")) {
3351 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3352 Name.starts_with(
"atomic.load.add.f64.p")) {
3355 Rep = Builder.CreateAtomicRMW(
3359 LLVMContext::MD_atomic_ignore_denormal_mode,
3364 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3365 Name.starts_with(
"atomic.load.dec.32.p")) {
3370 Rep = Builder.CreateAtomicRMW(
3374 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3380 Op.contains(
".cta.") ?
"block" :
"");
3381 if (
Op.starts_with(
"cas.")) {
3383 Value *Pair = Builder.CreateAtomicCmpXchg(
3386 Rep = Builder.CreateExtractValue(Pair, 0);
3404 "unexpected nvvm scoped atomic intrinsic");
3405 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3409 LLVMContext::MD_atomic_ignore_denormal_mode,
3413 }
else if (Name ==
"clz.ll") {
3416 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3417 {Arg, Builder.getFalse()},
3419 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3420 }
else if (Name ==
"popc.ll") {
3424 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3425 Arg,
nullptr,
"ctpop");
3426 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3427 }
else if (Name ==
"h2f") {
3429 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3430 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3431 }
else if (Name.consume_front(
"bitcast.") &&
3432 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3435 }
else if (Name ==
"rotate.b32") {
3438 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3439 {Arg, Arg, ShiftAmt});
3440 }
else if (Name ==
"rotate.b64") {
3444 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3445 {Arg, Arg, ZExtShiftAmt});
3446 }
else if (Name ==
"rotate.right.b64") {
3450 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3451 {Arg, Arg, ZExtShiftAmt});
3452 }
else if (Name ==
"swap.lo.hi.b64") {
3455 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3456 {Arg, Arg, Builder.getInt64(32)});
3457 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3460 Name.starts_with(
".to.gen"))) {
3462 }
else if (Name.consume_front(
"ldg.global")) {
3466 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3469 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3471 }
else if (Name ==
"tanh.approx.f32") {
3475 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3477 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3479 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3480 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3482 }
else if (Name ==
"barrier") {
3483 Rep = Builder.CreateIntrinsic(
3484 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3486 }
else if (Name ==
"barrier.sync") {
3487 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3489 }
else if (Name ==
"barrier.sync.cnt") {
3490 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3492 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3493 Name ==
"barrier0.or") {
3495 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3499 .
Case(
"barrier0.popc",
3500 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3501 .
Case(
"barrier0.and",
3502 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3503 .
Case(
"barrier0.or",
3504 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3505 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3506 Rep = Builder.CreateZExt(Bar, CI->
getType());
3520 ? Builder.CreateBitCast(Arg, NewType)
3523 Rep = Builder.CreateCall(NewFn, Args);
3524 if (
F->getReturnType()->isIntegerTy())
3525 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3535 Value *Rep =
nullptr;
3537 if (Name.starts_with(
"sse4a.movnt.")) {
3549 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3552 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3553 }
else if (Name.starts_with(
"avx.movnt.") ||
3554 Name.starts_with(
"avx512.storent.")) {
3566 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3567 }
else if (Name ==
"sse2.storel.dq") {
3572 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3573 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3574 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3575 }
else if (Name.starts_with(
"sse.storeu.") ||
3576 Name.starts_with(
"sse2.storeu.") ||
3577 Name.starts_with(
"avx.storeu.")) {
3580 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3581 }
else if (Name ==
"avx512.mask.store.ss") {
3585 }
else if (Name.starts_with(
"avx512.mask.store")) {
3587 bool Aligned = Name[17] !=
'u';
3590 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3593 bool CmpEq = Name[9] ==
'e';
3596 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3597 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3604 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3605 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3607 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3608 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3609 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3610 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3611 Name.starts_with(
"sse2.sqrt.p") ||
3612 Name.starts_with(
"sse.sqrt.p")) {
3613 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3614 {CI->getArgOperand(0)});
3615 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3619 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3620 : Intrinsic::x86_avx512_sqrt_pd_512;
3623 Rep = Builder.CreateIntrinsic(IID, Args);
3625 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3626 {CI->getArgOperand(0)});
3630 }
else if (Name.starts_with(
"avx512.ptestm") ||
3631 Name.starts_with(
"avx512.ptestnm")) {
3635 Rep = Builder.CreateAnd(Op0, Op1);
3641 Rep = Builder.CreateICmp(
Pred, Rep,
Zero);
3643 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3646 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3649 }
else if (Name.starts_with(
"avx512.kunpck")) {
3654 for (
unsigned i = 0; i != NumElts; ++i)
3663 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3664 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3665 }
else if (Name ==
"avx512.kand.w") {
3668 Rep = Builder.CreateAnd(
LHS,
RHS);
3669 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3670 }
else if (Name ==
"avx512.kandn.w") {
3673 LHS = Builder.CreateNot(
LHS);
3674 Rep = Builder.CreateAnd(
LHS,
RHS);
3675 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3676 }
else if (Name ==
"avx512.kor.w") {
3679 Rep = Builder.CreateOr(
LHS,
RHS);
3680 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3681 }
else if (Name ==
"avx512.kxor.w") {
3684 Rep = Builder.CreateXor(
LHS,
RHS);
3685 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3686 }
else if (Name ==
"avx512.kxnor.w") {
3689 LHS = Builder.CreateNot(
LHS);
3690 Rep = Builder.CreateXor(
LHS,
RHS);
3691 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3692 }
else if (Name ==
"avx512.knot.w") {
3694 Rep = Builder.CreateNot(Rep);
3695 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3696 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3699 Rep = Builder.CreateOr(
LHS,
RHS);
3700 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3702 if (Name[14] ==
'c')
3706 Rep = Builder.CreateICmpEQ(Rep,
C);
3707 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3708 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3709 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3710 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3711 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3714 ConstantInt::get(I32Ty, 0));
3716 ConstantInt::get(I32Ty, 0));
3718 if (Name.contains(
".add."))
3719 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3720 else if (Name.contains(
".sub."))
3721 EltOp = Builder.CreateFSub(Elt0, Elt1);
3722 else if (Name.contains(
".mul."))
3723 EltOp = Builder.CreateFMul(Elt0, Elt1);
3725 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3726 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3727 ConstantInt::get(I32Ty, 0));
3728 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3730 bool CmpEq = Name[16] ==
'e';
3732 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3734 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3741 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3744 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3747 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3754 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3756 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3757 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3759 if (VecWidth == 128 && EltWidth == 32)
3760 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3761 else if (VecWidth == 256 && EltWidth == 32)
3762 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3763 else if (VecWidth == 512 && EltWidth == 32)
3764 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3765 else if (VecWidth == 128 && EltWidth == 64)
3766 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3767 else if (VecWidth == 256 && EltWidth == 64)
3768 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3769 else if (VecWidth == 512 && EltWidth == 64)
3770 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3777 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3780 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3781 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3783 if (VecWidth == 128 && EltWidth == 32)
3784 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3785 else if (VecWidth == 256 && EltWidth == 32)
3786 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3787 else if (VecWidth == 512 && EltWidth == 32)
3788 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3789 else if (VecWidth == 128 && EltWidth == 64)
3790 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3791 else if (VecWidth == 256 && EltWidth == 64)
3792 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3793 else if (VecWidth == 512 && EltWidth == 64)
3794 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3799 if (VecWidth == 512)
3801 Args.push_back(Mask);
3803 Rep = Builder.CreateIntrinsic(IID, Args);
3804 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3808 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3811 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3812 Name.starts_with(
"avx512.cvtw2mask.") ||
3813 Name.starts_with(
"avx512.cvtd2mask.") ||
3814 Name.starts_with(
"avx512.cvtq2mask.")) {
3819 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3820 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3821 Name.starts_with(
"avx512.mask.pabs")) {
3823 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3824 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3825 Name.starts_with(
"avx512.mask.pmaxs")) {
3827 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3828 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3829 Name.starts_with(
"avx512.mask.pmaxu")) {
3831 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3832 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3833 Name.starts_with(
"avx512.mask.pmins")) {
3835 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3836 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3837 Name.starts_with(
"avx512.mask.pminu")) {
3839 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3840 Name.starts_with(
"avx512.pmulh.w")) {
3842 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3843 Name.starts_with(
"avx512.pmulhu.w")) {
3845 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3846 Name ==
"avx512.pmulu.dq.512" ||
3847 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3849 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3850 Name ==
"avx512.pmul.dq.512" ||
3851 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3853 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3854 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3859 }
else if (Name ==
"avx512.cvtusi2sd") {
3864 }
else if (Name ==
"sse2.cvtss2sd") {
3866 Rep = Builder.CreateFPExt(
3869 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3870 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3871 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3872 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3873 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3874 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3875 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3876 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3877 Name ==
"avx512.mask.cvtqq2ps.256" ||
3878 Name ==
"avx512.mask.cvtqq2ps.512" ||
3879 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3880 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3881 Name ==
"avx.cvt.ps2.pd.256" ||
3882 Name ==
"avx512.mask.cvtps2pd.128" ||
3883 Name ==
"avx512.mask.cvtps2pd.256") {
3888 unsigned NumDstElts = DstTy->getNumElements();
3889 if (NumDstElts < SrcTy->getNumElements()) {
3890 assert(NumDstElts == 2 &&
"Unexpected vector size");
3891 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3894 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3895 bool IsUnsigned = Name.contains(
"cvtu");
3897 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3901 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3902 : Intrinsic::x86_avx512_sitofp_round;
3903 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3906 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3907 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3913 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3914 Name.starts_with(
"vcvtph2ps.")) {
3918 unsigned NumDstElts = DstTy->getNumElements();
3919 if (NumDstElts != SrcTy->getNumElements()) {
3920 assert(NumDstElts == 4 &&
"Unexpected vector size");
3921 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3923 Rep = Builder.CreateBitCast(
3925 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3929 }
else if (Name.starts_with(
"avx512.mask.load")) {
3931 bool Aligned = Name[16] !=
'u';
3934 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3938 ResultTy->getNumElements());
3939 Rep = Builder.CreateIntrinsic(
3940 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3942 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3948 Rep = Builder.CreateIntrinsic(
3949 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3951 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3952 Name.starts_with(
"avx512.mask.expand.")) {
3956 ResultTy->getNumElements());
3958 bool IsCompress = Name[12] ==
'c';
3959 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3960 : Intrinsic::x86_avx512_mask_expand;
3961 Rep = Builder.CreateIntrinsic(
3963 }
else if (Name.starts_with(
"xop.vpcom")) {
3965 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3966 Name.ends_with(
"uq"))
3968 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3969 Name.ends_with(
"d") || Name.ends_with(
"q"))
3978 Name = Name.substr(9);
3979 if (Name.starts_with(
"lt"))
3981 else if (Name.starts_with(
"le"))
3983 else if (Name.starts_with(
"gt"))
3985 else if (Name.starts_with(
"ge"))
3987 else if (Name.starts_with(
"eq"))
3989 else if (Name.starts_with(
"ne"))
3991 else if (Name.starts_with(
"false"))
3993 else if (Name.starts_with(
"true"))
4000 }
else if (Name.starts_with(
"xop.vpcmov")) {
4002 Value *NotSel = Builder.CreateNot(Sel);
4005 Rep = Builder.CreateOr(Sel0, Sel1);
4006 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
4007 Name.starts_with(
"avx512.mask.prol")) {
4009 }
else if (Name.starts_with(
"avx512.pror") ||
4010 Name.starts_with(
"avx512.mask.pror")) {
4012 }
else if (Name.starts_with(
"avx512.vpshld.") ||
4013 Name.starts_with(
"avx512.mask.vpshld") ||
4014 Name.starts_with(
"avx512.maskz.vpshld")) {
4015 bool ZeroMask = Name[11] ==
'z';
4017 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
4018 Name.starts_with(
"avx512.mask.vpshrd") ||
4019 Name.starts_with(
"avx512.maskz.vpshrd")) {
4020 bool ZeroMask = Name[11] ==
'z';
4022 }
else if (Name ==
"sse42.crc32.64.8") {
4025 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
4027 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
4028 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
4029 Name.starts_with(
"avx512.vbroadcast.s")) {
4032 Type *EltTy = VecTy->getElementType();
4033 unsigned EltNum = VecTy->getNumElements();
4037 for (
unsigned I = 0;
I < EltNum; ++
I)
4038 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
4039 }
else if (Name.starts_with(
"sse41.pmovsx") ||
4040 Name.starts_with(
"sse41.pmovzx") ||
4041 Name.starts_with(
"avx2.pmovsx") ||
4042 Name.starts_with(
"avx2.pmovzx") ||
4043 Name.starts_with(
"avx512.mask.pmovsx") ||
4044 Name.starts_with(
"avx512.mask.pmovzx")) {
4046 unsigned NumDstElts = DstTy->getNumElements();
4050 for (
unsigned i = 0; i != NumDstElts; ++i)
4055 bool DoSext = Name.contains(
"pmovsx");
4057 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
4062 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
4063 Name ==
"avx512.mask.pmov.qd.512" ||
4064 Name ==
"avx512.mask.pmov.wb.256" ||
4065 Name ==
"avx512.mask.pmov.wb.512") {
4070 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4071 Name ==
"avx2.vbroadcasti128") {
4077 if (NumSrcElts == 2)
4080 Rep = Builder.CreateShuffleVector(
Load,
4082 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4083 Name.starts_with(
"avx512.mask.shuf.f")) {
4088 unsigned ControlBitsMask = NumLanes - 1;
4089 unsigned NumControlBits = NumLanes / 2;
4092 for (
unsigned l = 0; l != NumLanes; ++l) {
4093 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4095 if (l >= NumLanes / 2)
4096 LaneMask += NumLanes;
4097 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4098 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4104 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4105 Name.starts_with(
"avx512.mask.broadcasti")) {
4108 unsigned NumDstElts =
4112 for (
unsigned i = 0; i != NumDstElts; ++i)
4113 ShuffleMask[i] = i % NumSrcElts;
4119 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4120 Name.starts_with(
"avx2.vbroadcast") ||
4121 Name.starts_with(
"avx512.pbroadcast") ||
4122 Name.starts_with(
"avx512.mask.broadcast.s")) {
4129 Rep = Builder.CreateShuffleVector(
Op, M);
4134 }
else if (Name.starts_with(
"sse2.padds.") ||
4135 Name.starts_with(
"avx2.padds.") ||
4136 Name.starts_with(
"avx512.padds.") ||
4137 Name.starts_with(
"avx512.mask.padds.")) {
4139 }
else if (Name.starts_with(
"sse2.psubs.") ||
4140 Name.starts_with(
"avx2.psubs.") ||
4141 Name.starts_with(
"avx512.psubs.") ||
4142 Name.starts_with(
"avx512.mask.psubs.")) {
4144 }
else if (Name.starts_with(
"sse2.paddus.") ||
4145 Name.starts_with(
"avx2.paddus.") ||
4146 Name.starts_with(
"avx512.mask.paddus.")) {
4148 }
else if (Name.starts_with(
"sse2.psubus.") ||
4149 Name.starts_with(
"avx2.psubus.") ||
4150 Name.starts_with(
"avx512.mask.psubus.")) {
4152 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4157 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4161 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4166 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4171 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4172 Name ==
"avx512.psll.dq.512") {
4176 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4177 Name ==
"avx512.psrl.dq.512") {
4181 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4182 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4183 Name.starts_with(
"avx2.pblendd.")) {
4188 unsigned NumElts = VecTy->getNumElements();
4191 for (
unsigned i = 0; i != NumElts; ++i)
4192 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4194 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4195 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4196 Name ==
"avx2.vinserti128" ||
4197 Name.starts_with(
"avx512.mask.insert")) {
4201 unsigned DstNumElts =
4203 unsigned SrcNumElts =
4205 unsigned Scale = DstNumElts / SrcNumElts;
4212 for (
unsigned i = 0; i != SrcNumElts; ++i)
4214 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4215 Idxs[i] = SrcNumElts;
4216 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4230 for (
unsigned i = 0; i != DstNumElts; ++i)
4233 for (
unsigned i = 0; i != SrcNumElts; ++i)
4234 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4235 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4241 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4242 Name ==
"avx2.vextracti128" ||
4243 Name.starts_with(
"avx512.mask.vextract")) {
4246 unsigned DstNumElts =
4248 unsigned SrcNumElts =
4250 unsigned Scale = SrcNumElts / DstNumElts;
4257 for (
unsigned i = 0; i != DstNumElts; ++i) {
4258 Idxs[i] = i + (
Imm * DstNumElts);
4260 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4266 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4267 Name.starts_with(
"avx512.mask.perm.di.")) {
4271 unsigned NumElts = VecTy->getNumElements();
4274 for (
unsigned i = 0; i != NumElts; ++i)
4275 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4277 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4282 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4294 unsigned HalfSize = NumElts / 2;
4306 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4307 for (
unsigned i = 0; i < HalfSize; ++i)
4308 ShuffleMask[i] = StartIndex + i;
4311 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4312 for (
unsigned i = 0; i < HalfSize; ++i)
4313 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4315 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4317 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4318 Name.starts_with(
"avx512.mask.vpermil.p") ||
4319 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4323 unsigned NumElts = VecTy->getNumElements();
4325 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4326 unsigned IdxMask = ((1 << IdxSize) - 1);
4332 for (
unsigned i = 0; i != NumElts; ++i)
4333 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4335 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4340 }
else if (Name ==
"sse2.pshufl.w" ||
4341 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4346 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4350 for (
unsigned l = 0; l != NumElts; l += 8) {
4351 for (
unsigned i = 0; i != 4; ++i)
4352 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4353 for (
unsigned i = 4; i != 8; ++i)
4354 Idxs[i + l] = i + l;
4357 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4362 }
else if (Name ==
"sse2.pshufh.w" ||
4363 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4368 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4372 for (
unsigned l = 0; l != NumElts; l += 8) {
4373 for (
unsigned i = 0; i != 4; ++i)
4374 Idxs[i + l] = i + l;
4375 for (
unsigned i = 0; i != 4; ++i)
4376 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4379 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4384 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4391 unsigned HalfLaneElts = NumLaneElts / 2;
4394 for (
unsigned i = 0; i != NumElts; ++i) {
4396 Idxs[i] = i - (i % NumLaneElts);
4398 if ((i % NumLaneElts) >= HalfLaneElts)
4402 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4405 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4409 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4410 Name.starts_with(
"avx512.mask.movshdup") ||
4411 Name.starts_with(
"avx512.mask.movsldup")) {
4417 if (Name.starts_with(
"avx512.mask.movshdup."))
4421 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4422 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4423 Idxs[i + l + 0] = i + l +
Offset;
4424 Idxs[i + l + 1] = i + l +
Offset;
4427 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4431 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4432 Name.starts_with(
"avx512.mask.unpckl.")) {
4439 for (
int l = 0; l != NumElts; l += NumLaneElts)
4440 for (
int i = 0; i != NumLaneElts; ++i)
4441 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4443 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4447 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4448 Name.starts_with(
"avx512.mask.unpckh.")) {
4455 for (
int l = 0; l != NumElts; l += NumLaneElts)
4456 for (
int i = 0; i != NumLaneElts; ++i)
4457 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4459 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4463 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4464 Name.starts_with(
"avx512.mask.pand.")) {
4467 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4469 Rep = Builder.CreateBitCast(Rep, FTy);
4472 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4473 Name.starts_with(
"avx512.mask.pandn.")) {
4476 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4477 Rep = Builder.CreateAnd(Rep,
4479 Rep = Builder.CreateBitCast(Rep, FTy);
4482 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4483 Name.starts_with(
"avx512.mask.por.")) {
4486 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4488 Rep = Builder.CreateBitCast(Rep, FTy);
4491 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4492 Name.starts_with(
"avx512.mask.pxor.")) {
4495 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4497 Rep = Builder.CreateBitCast(Rep, FTy);
4500 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4504 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4508 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4512 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4513 if (Name.ends_with(
".512")) {
4515 if (Name[17] ==
's')
4516 IID = Intrinsic::x86_avx512_add_ps_512;
4518 IID = Intrinsic::x86_avx512_add_pd_512;
4520 Rep = Builder.CreateIntrinsic(
4528 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4529 if (Name.ends_with(
".512")) {
4531 if (Name[17] ==
's')
4532 IID = Intrinsic::x86_avx512_div_ps_512;
4534 IID = Intrinsic::x86_avx512_div_pd_512;
4536 Rep = Builder.CreateIntrinsic(
4544 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4545 if (Name.ends_with(
".512")) {
4547 if (Name[17] ==
's')
4548 IID = Intrinsic::x86_avx512_mul_ps_512;
4550 IID = Intrinsic::x86_avx512_mul_pd_512;
4552 Rep = Builder.CreateIntrinsic(
4560 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4561 if (Name.ends_with(
".512")) {
4563 if (Name[17] ==
's')
4564 IID = Intrinsic::x86_avx512_sub_ps_512;
4566 IID = Intrinsic::x86_avx512_sub_pd_512;
4568 Rep = Builder.CreateIntrinsic(
4576 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4577 Name.starts_with(
"avx512.mask.min.p")) &&
4578 Name.drop_front(18) ==
".512") {
4579 bool IsDouble = Name[17] ==
'd';
4580 bool IsMin = Name[13] ==
'i';
4582 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4583 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4586 Rep = Builder.CreateIntrinsic(
4591 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4593 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4594 {CI->getArgOperand(0), Builder.getInt1(false)});
4597 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4598 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4599 bool IsVariable = Name[16] ==
'v';
4600 char Size = Name[16] ==
'.' ? Name[17]
4601 : Name[17] ==
'.' ? Name[18]
4602 : Name[18] ==
'.' ? Name[19]
4606 if (IsVariable && Name[17] !=
'.') {
4607 if (
Size ==
'd' && Name[17] ==
'2')
4608 IID = Intrinsic::x86_avx2_psllv_q;
4609 else if (
Size ==
'd' && Name[17] ==
'4')
4610 IID = Intrinsic::x86_avx2_psllv_q_256;
4611 else if (
Size ==
's' && Name[17] ==
'4')
4612 IID = Intrinsic::x86_avx2_psllv_d;
4613 else if (
Size ==
's' && Name[17] ==
'8')
4614 IID = Intrinsic::x86_avx2_psllv_d_256;
4615 else if (
Size ==
'h' && Name[17] ==
'8')
4616 IID = Intrinsic::x86_avx512_psllv_w_128;
4617 else if (
Size ==
'h' && Name[17] ==
'1')
4618 IID = Intrinsic::x86_avx512_psllv_w_256;
4619 else if (Name[17] ==
'3' && Name[18] ==
'2')
4620 IID = Intrinsic::x86_avx512_psllv_w_512;
4623 }
else if (Name.ends_with(
".128")) {
4625 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4626 : Intrinsic::x86_sse2_psll_d;
4627 else if (
Size ==
'q')
4628 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4629 : Intrinsic::x86_sse2_psll_q;
4630 else if (
Size ==
'w')
4631 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4632 : Intrinsic::x86_sse2_psll_w;
4635 }
else if (Name.ends_with(
".256")) {
4637 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4638 : Intrinsic::x86_avx2_psll_d;
4639 else if (
Size ==
'q')
4640 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4641 : Intrinsic::x86_avx2_psll_q;
4642 else if (
Size ==
'w')
4643 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4644 : Intrinsic::x86_avx2_psll_w;
4649 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4650 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4651 : Intrinsic::x86_avx512_psll_d_512;
4652 else if (
Size ==
'q')
4653 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4654 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4655 : Intrinsic::x86_avx512_psll_q_512;
4656 else if (
Size ==
'w')
4657 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4658 : Intrinsic::x86_avx512_psll_w_512;
4664 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4665 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4666 bool IsVariable = Name[16] ==
'v';
4667 char Size = Name[16] ==
'.' ? Name[17]
4668 : Name[17] ==
'.' ? Name[18]
4669 : Name[18] ==
'.' ? Name[19]
4673 if (IsVariable && Name[17] !=
'.') {
4674 if (
Size ==
'd' && Name[17] ==
'2')
4675 IID = Intrinsic::x86_avx2_psrlv_q;
4676 else if (
Size ==
'd' && Name[17] ==
'4')
4677 IID = Intrinsic::x86_avx2_psrlv_q_256;
4678 else if (
Size ==
's' && Name[17] ==
'4')
4679 IID = Intrinsic::x86_avx2_psrlv_d;
4680 else if (
Size ==
's' && Name[17] ==
'8')
4681 IID = Intrinsic::x86_avx2_psrlv_d_256;
4682 else if (
Size ==
'h' && Name[17] ==
'8')
4683 IID = Intrinsic::x86_avx512_psrlv_w_128;
4684 else if (
Size ==
'h' && Name[17] ==
'1')
4685 IID = Intrinsic::x86_avx512_psrlv_w_256;
4686 else if (Name[17] ==
'3' && Name[18] ==
'2')
4687 IID = Intrinsic::x86_avx512_psrlv_w_512;
4690 }
else if (Name.ends_with(
".128")) {
4692 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4693 : Intrinsic::x86_sse2_psrl_d;
4694 else if (
Size ==
'q')
4695 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4696 : Intrinsic::x86_sse2_psrl_q;
4697 else if (
Size ==
'w')
4698 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4699 : Intrinsic::x86_sse2_psrl_w;
4702 }
else if (Name.ends_with(
".256")) {
4704 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4705 : Intrinsic::x86_avx2_psrl_d;
4706 else if (
Size ==
'q')
4707 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4708 : Intrinsic::x86_avx2_psrl_q;
4709 else if (
Size ==
'w')
4710 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4711 : Intrinsic::x86_avx2_psrl_w;
4716 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4717 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4718 : Intrinsic::x86_avx512_psrl_d_512;
4719 else if (
Size ==
'q')
4720 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4721 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4722 : Intrinsic::x86_avx512_psrl_q_512;
4723 else if (
Size ==
'w')
4724 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4725 : Intrinsic::x86_avx512_psrl_w_512;
4731 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4732 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4733 bool IsVariable = Name[16] ==
'v';
4734 char Size = Name[16] ==
'.' ? Name[17]
4735 : Name[17] ==
'.' ? Name[18]
4736 : Name[18] ==
'.' ? Name[19]
4740 if (IsVariable && Name[17] !=
'.') {
4741 if (
Size ==
's' && Name[17] ==
'4')
4742 IID = Intrinsic::x86_avx2_psrav_d;
4743 else if (
Size ==
's' && Name[17] ==
'8')
4744 IID = Intrinsic::x86_avx2_psrav_d_256;
4745 else if (
Size ==
'h' && Name[17] ==
'8')
4746 IID = Intrinsic::x86_avx512_psrav_w_128;
4747 else if (
Size ==
'h' && Name[17] ==
'1')
4748 IID = Intrinsic::x86_avx512_psrav_w_256;
4749 else if (Name[17] ==
'3' && Name[18] ==
'2')
4750 IID = Intrinsic::x86_avx512_psrav_w_512;
4753 }
else if (Name.ends_with(
".128")) {
4755 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4756 : Intrinsic::x86_sse2_psra_d;
4757 else if (
Size ==
'q')
4758 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4759 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4760 : Intrinsic::x86_avx512_psra_q_128;
4761 else if (
Size ==
'w')
4762 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4763 : Intrinsic::x86_sse2_psra_w;
4766 }
else if (Name.ends_with(
".256")) {
4768 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4769 : Intrinsic::x86_avx2_psra_d;
4770 else if (
Size ==
'q')
4771 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4772 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4773 : Intrinsic::x86_avx512_psra_q_256;
4774 else if (
Size ==
'w')
4775 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4776 : Intrinsic::x86_avx2_psra_w;
4781 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4782 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4783 : Intrinsic::x86_avx512_psra_d_512;
4784 else if (
Size ==
'q')
4785 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4786 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4787 : Intrinsic::x86_avx512_psra_q_512;
4788 else if (
Size ==
'w')
4789 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4790 : Intrinsic::x86_avx512_psra_w_512;
4796 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4798 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4800 }
else if (Name.ends_with(
".movntdqa")) {
4804 LoadInst *LI = Builder.CreateAlignedLoad(
4809 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4810 Name.starts_with(
"fma.vfmsub.") ||
4811 Name.starts_with(
"fma.vfnmadd.") ||
4812 Name.starts_with(
"fma.vfnmsub.")) {
4813 bool NegMul = Name[6] ==
'n';
4814 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4815 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4826 if (NegMul && !IsScalar)
4827 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4828 if (NegMul && IsScalar)
4829 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4831 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4833 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4837 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4845 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4849 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4850 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4851 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4852 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4853 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4854 bool IsMask3 = Name[11] ==
'3';
4855 bool IsMaskZ = Name[11] ==
'z';
4857 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4858 bool NegMul = Name[2] ==
'n';
4859 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4865 if (NegMul && (IsMask3 || IsMaskZ))
4866 A = Builder.CreateFNeg(
A);
4867 if (NegMul && !(IsMask3 || IsMaskZ))
4868 B = Builder.CreateFNeg(
B);
4870 C = Builder.CreateFNeg(
C);
4872 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4873 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4874 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4881 if (Name.back() ==
'd')
4882 IID = Intrinsic::x86_avx512_vfmadd_f64;
4884 IID = Intrinsic::x86_avx512_vfmadd_f32;
4885 Rep = Builder.CreateIntrinsic(IID,
Ops);
4887 Rep = Builder.CreateFMA(
A,
B,
C);
4896 if (NegAcc && IsMask3)
4901 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4903 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4904 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4905 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4906 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4907 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4908 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4909 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4910 bool IsMask3 = Name[11] ==
'3';
4911 bool IsMaskZ = Name[11] ==
'z';
4913 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4914 bool NegMul = Name[2] ==
'n';
4915 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4921 if (NegMul && (IsMask3 || IsMaskZ))
4922 A = Builder.CreateFNeg(
A);
4923 if (NegMul && !(IsMask3 || IsMaskZ))
4924 B = Builder.CreateFNeg(
B);
4926 C = Builder.CreateFNeg(
C);
4933 if (Name[Name.size() - 5] ==
's')
4934 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4936 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4940 Rep = Builder.CreateFMA(
A,
B,
C);
4948 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4952 if (VecWidth == 128 && EltWidth == 32)
4953 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4954 else if (VecWidth == 256 && EltWidth == 32)
4955 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4956 else if (VecWidth == 128 && EltWidth == 64)
4957 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4958 else if (VecWidth == 256 && EltWidth == 64)
4959 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4965 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4966 Rep = Builder.CreateIntrinsic(IID,
Ops);
4967 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4968 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4969 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4970 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4971 bool IsMask3 = Name[11] ==
'3';
4972 bool IsMaskZ = Name[11] ==
'z';
4974 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4975 bool IsSubAdd = Name[3] ==
's';
4979 if (Name[Name.size() - 5] ==
's')
4980 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4982 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4987 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4989 Rep = Builder.CreateIntrinsic(IID,
Ops);
4998 Value *Odd = Builder.CreateCall(FMA,
Ops);
4999 Ops[2] = Builder.CreateFNeg(
Ops[2]);
5000 Value *Even = Builder.CreateCall(FMA,
Ops);
5006 for (
int i = 0; i != NumElts; ++i)
5007 Idxs[i] = i + (i % 2) * NumElts;
5009 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
5017 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
5018 Name.starts_with(
"avx512.maskz.pternlog.")) {
5019 bool ZeroMask = Name[11] ==
'z';
5023 if (VecWidth == 128 && EltWidth == 32)
5024 IID = Intrinsic::x86_avx512_pternlog_d_128;
5025 else if (VecWidth == 256 && EltWidth == 32)
5026 IID = Intrinsic::x86_avx512_pternlog_d_256;
5027 else if (VecWidth == 512 && EltWidth == 32)
5028 IID = Intrinsic::x86_avx512_pternlog_d_512;
5029 else if (VecWidth == 128 && EltWidth == 64)
5030 IID = Intrinsic::x86_avx512_pternlog_q_128;
5031 else if (VecWidth == 256 && EltWidth == 64)
5032 IID = Intrinsic::x86_avx512_pternlog_q_256;
5033 else if (VecWidth == 512 && EltWidth == 64)
5034 IID = Intrinsic::x86_avx512_pternlog_q_512;
5040 Rep = Builder.CreateIntrinsic(IID, Args);
5044 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
5045 Name.starts_with(
"avx512.maskz.vpmadd52")) {
5046 bool ZeroMask = Name[11] ==
'z';
5047 bool High = Name[20] ==
'h' || Name[21] ==
'h';
5050 if (VecWidth == 128 && !
High)
5051 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
5052 else if (VecWidth == 256 && !
High)
5053 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
5054 else if (VecWidth == 512 && !
High)
5055 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
5056 else if (VecWidth == 128 &&
High)
5057 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
5058 else if (VecWidth == 256 &&
High)
5059 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
5060 else if (VecWidth == 512 &&
High)
5061 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5067 Rep = Builder.CreateIntrinsic(IID, Args);
5071 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5072 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5073 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5074 bool ZeroMask = Name[11] ==
'z';
5075 bool IndexForm = Name[17] ==
'i';
5077 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5078 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5079 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5080 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5081 bool ZeroMask = Name[11] ==
'z';
5082 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5085 if (VecWidth == 128 && !IsSaturating)
5086 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5087 else if (VecWidth == 256 && !IsSaturating)
5088 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5089 else if (VecWidth == 512 && !IsSaturating)
5090 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5091 else if (VecWidth == 128 && IsSaturating)
5092 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5093 else if (VecWidth == 256 && IsSaturating)
5094 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5095 else if (VecWidth == 512 && IsSaturating)
5096 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5106 if (Args[1]->
getType()->isVectorTy() &&
5109 ->isIntegerTy(32) &&
5110 Args[2]->
getType()->isVectorTy() &&
5113 ->isIntegerTy(32)) {
5114 Type *NewArgType =
nullptr;
5115 if (VecWidth == 128)
5117 else if (VecWidth == 256)
5119 else if (VecWidth == 512)
5125 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5126 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5129 Rep = Builder.CreateIntrinsic(IID, Args);
5133 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5134 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5135 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5136 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5137 bool ZeroMask = Name[11] ==
'z';
5138 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5141 if (VecWidth == 128 && !IsSaturating)
5142 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5143 else if (VecWidth == 256 && !IsSaturating)
5144 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5145 else if (VecWidth == 512 && !IsSaturating)
5146 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5147 else if (VecWidth == 128 && IsSaturating)
5148 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5149 else if (VecWidth == 256 && IsSaturating)
5150 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5151 else if (VecWidth == 512 && IsSaturating)
5152 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5162 if (Args[1]->
getType()->isVectorTy() &&
5165 ->isIntegerTy(32) &&
5166 Args[2]->
getType()->isVectorTy() &&
5169 ->isIntegerTy(32)) {
5170 Type *NewArgType =
nullptr;
5171 if (VecWidth == 128)
5173 else if (VecWidth == 256)
5175 else if (VecWidth == 512)
5181 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5182 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5185 Rep = Builder.CreateIntrinsic(IID, Args);
5189 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5190 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5191 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5193 if (Name[0] ==
'a' && Name.back() ==
'2')
5194 IID = Intrinsic::x86_addcarry_32;
5195 else if (Name[0] ==
'a' && Name.back() ==
'4')
5196 IID = Intrinsic::x86_addcarry_64;
5197 else if (Name[0] ==
's' && Name.back() ==
'2')
5198 IID = Intrinsic::x86_subborrow_32;
5199 else if (Name[0] ==
's' && Name.back() ==
'4')
5200 IID = Intrinsic::x86_subborrow_64;
5207 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5210 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5213 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5217 }
else if (Name.starts_with(
"avx512.mask.") &&
5220 }
else if (Name.starts_with(
"bmi.pdep.")) {
5222 }
else if (Name.starts_with(
"bmi.pext.")) {
5232 if (Name.starts_with(
"neon.bfcvt")) {
5233 if (Name.starts_with(
"neon.bfcvtn2")) {
5235 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5237 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5238 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5241 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5242 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5244 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5248 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5249 return Builder.CreateShuffleVector(
5252 return Builder.CreateFPTrunc(CI->
getOperand(0),
5255 }
else if (Name.starts_with(
"sve.fcvt")) {
5258 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5259 .
Case(
"sve.fcvtnt.bf16f32",
5260 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5272 if (Args[1]->
getType() != BadPredTy)
5275 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5276 BadPredTy, Args[1]);
5277 Args[1] = Builder.CreateIntrinsic(
5278 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5280 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5284 if (Name ==
"neon.vcvtfp2hf")
5285 return Builder.CreateBitCast(
5286 Builder.CreateFPTrunc(
5290 if (Name ==
"neon.vcvthf2fp")
5291 return Builder.CreateFPExt(
5292 Builder.CreateBitCast(
5302 if (Name ==
"mve.vctp64.old") {
5305 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5308 Value *
C1 = Builder.CreateIntrinsic(
5309 Intrinsic::arm_mve_pred_v2i,
5311 return Builder.CreateIntrinsic(
5312 Intrinsic::arm_mve_pred_i2v,
5314 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5315 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5316 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5317 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5319 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5320 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5321 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5322 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5324 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5325 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5326 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5327 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5328 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5329 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5330 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5331 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5332 std::vector<Type *> Tys;
5336 case Intrinsic::arm_mve_mull_int_predicated:
5337 case Intrinsic::arm_mve_vqdmull_predicated:
5338 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5341 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5342 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5343 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5347 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5351 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5355 case Intrinsic::arm_cde_vcx1q_predicated:
5356 case Intrinsic::arm_cde_vcx1qa_predicated:
5357 case Intrinsic::arm_cde_vcx2q_predicated:
5358 case Intrinsic::arm_cde_vcx2qa_predicated:
5359 case Intrinsic::arm_cde_vcx3q_predicated:
5360 case Intrinsic::arm_cde_vcx3qa_predicated:
5367 std::vector<Value *>
Ops;
5369 Type *Ty =
Op->getType();
5370 if (Ty->getScalarSizeInBits() == 1) {
5371 Value *
C1 = Builder.CreateIntrinsic(
5372 Intrinsic::arm_mve_pred_v2i,
5374 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty},
C1);
5379 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5394 auto UpgradeLegacyWMMAIUIntrinsicCall =
5399 Args.push_back(Builder.getFalse());
5403 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5410 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5414 NewCall->copyMetadata(*CI);
5418 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5419 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5420 "intrinsic should have 7 arguments");
5423 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5425 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5426 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5427 "intrinsic should have 8 arguments");
5432 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5435 switch (
F->getIntrinsicID()) {
5438 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5439 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5440 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5441 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5442 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5443 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5458 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5461 F->getParent(),
F->getIntrinsicID(), Overloads);
5466 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5470 NewCall->copyMetadata(*CI);
5471 NewCall->takeName(CI);
5476 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5482 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5483 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5491 if (Name.starts_with(
"addrspacecast.nonnull")) {
5494 Value *ASC = Builder.CreateAddrSpaceCast(
5517 if (NumOperands < 3)
5530 bool IsVolatile =
false;
5534 if (NumOperands > 3)
5539 if (NumOperands > 5) {
5541 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5555 if (VT->getElementType()->isIntegerTy(16)) {
5558 Val = Builder.CreateBitCast(Val, AsBF16);
5566 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5568 unsigned AddrSpace = PtrTy->getAddressSpace();
5571 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5573 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5578 MDNode *RangeNotPrivate =
5581 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5587 return Builder.CreateBitCast(RMW, RetTy);
5608 return MAV->getMetadata();
5617 if (Name ==
"label") {
5619 }
else if (Name ==
"assign") {
5626 }
else if (Name ==
"declare") {
5630 }
else if (Name ==
"addr") {
5640 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5641 }
else if (Name ==
"value") {
5644 unsigned ExprOp = 2;
5659 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5667 int64_t OffsetVal =
Offset->getSExtValue();
5668 return Builder.CreateIntrinsic(OffsetVal >= 0
5669 ? Intrinsic::vector_splice_left
5670 : Intrinsic::vector_splice_right,
5672 {CI->getArgOperand(0), CI->getArgOperand(1),
5673 Builder.getInt32(std::abs(OffsetVal))});
5678 if (Name.starts_with(
"to.fp16")) {
5680 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5681 return Builder.CreateBitCast(Cast, CI->
getType());
5684 if (Name.starts_with(
"from.fp16")) {
5686 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5687 return Builder.CreateFPExt(Cast, CI->
getType());
5746 else if (Opcode == Instruction::ICmp)
5749 else if (Opcode == Instruction::FCmp)
5752 else if (Opcode == Instruction::Select)
5757 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5769 if (Defaults.empty())
5772 unsigned OldArgCount = CI->
arg_size();
5773 unsigned NewArgCount = NewFn->
arg_size();
5775 if (OldArgCount < FirstDefault)
5779 if (OldArgCount > NewArgCount)
5784 if (OldArgCount == NewArgCount) {
5796 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5797 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5798 "missing argument outside the default range");
5799 Type *ParamTy = NewFT->getParamType(Idx);
5804 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5810 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5841 if (!Name.consume_front(
"llvm."))
5844 bool IsX86 = Name.consume_front(
"x86.");
5845 bool IsNVVM = Name.consume_front(
"nvvm.");
5846 bool IsAArch64 = Name.consume_front(
"aarch64.");
5847 bool IsARM = Name.consume_front(
"arm.");
5848 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5849 bool IsDbg = Name.consume_front(
"dbg.");
5851 (Name.consume_front(
"experimental.vector.splice") ||
5852 Name.consume_front(
"vector.splice")) &&
5853 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5854 Value *Rep =
nullptr;
5856 if (!IsX86 && Name ==
"stackprotectorcheck") {
5858 }
else if (IsNVVM) {
5862 }
else if (IsAArch64) {
5866 }
else if (IsAMDGCN) {
5870 }
else if (IsOldSplice) {
5872 }
else if (Name.consume_front(
"convert.")) {
5874 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5889 const auto &DefaultCase = [&]() ->
void {
5897 "Unknown function for CallBase upgrade and isn't just a name change");
5905 "Return type must have changed");
5906 assert(OldST->getNumElements() ==
5908 "Must have same number of elements");
5911 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5914 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5915 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5916 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5937 case Intrinsic::arm_neon_vst1:
5938 case Intrinsic::arm_neon_vst2:
5939 case Intrinsic::arm_neon_vst3:
5940 case Intrinsic::arm_neon_vst4:
5941 case Intrinsic::arm_neon_vst2lane:
5942 case Intrinsic::arm_neon_vst3lane:
5943 case Intrinsic::arm_neon_vst4lane: {
5945 NewCall = Builder.CreateCall(NewFn, Args);
5948 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5949 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5950 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5955 NewCall = Builder.CreateCall(NewFn, Args);
5958 case Intrinsic::aarch64_sve_ld3_sret:
5959 case Intrinsic::aarch64_sve_ld4_sret:
5960 case Intrinsic::aarch64_sve_ld2_sret: {
5968 Name = Name.substr(5);
5975 unsigned MinElts = RetTy->getMinNumElements() /
N;
5977 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5979 for (
unsigned I = 0;
I <
N;
I++) {
5980 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5981 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5987 case Intrinsic::coro_end_async:
5988 case Intrinsic::coro_end: {
5990 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5992 NewCall = Builder.CreateCall(NewFn, Args);
5997 CI->
getModule(), Intrinsic::coro_is_in_ramp);
5998 Value *InRamp = Builder.CreateCall(IsInRamp);
6008 case Intrinsic::vector_extract: {
6010 Name = Name.substr(5);
6011 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
6016 unsigned MinElts = RetTy->getMinNumElements();
6019 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
6023 case Intrinsic::vector_insert: {
6025 Name = Name.substr(5);
6026 if (!Name.starts_with(
"aarch64.sve.tuple")) {
6030 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
6035 NewCall = Builder.CreateCall(
6039 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
6045 assert(
N > 1 &&
"Create is expected to be between 2-4");
6048 unsigned MinElts = RetTy->getMinNumElements() /
N;
6049 for (
unsigned I = 0;
I <
N;
I++) {
6051 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
6058 case Intrinsic::arm_neon_bfdot:
6059 case Intrinsic::arm_neon_bfmmla:
6060 case Intrinsic::arm_neon_bfmlalb:
6061 case Intrinsic::arm_neon_bfmlalt:
6062 case Intrinsic::aarch64_neon_bfdot:
6063 case Intrinsic::aarch64_neon_bfmmla:
6064 case Intrinsic::aarch64_neon_bfmlalb:
6065 case Intrinsic::aarch64_neon_bfmlalt: {
6068 "Mismatch between function args and call args");
6069 size_t OperandWidth =
6071 assert((OperandWidth == 64 || OperandWidth == 128) &&
6072 "Unexpected operand width");
6074 auto Iter = CI->
args().begin();
6075 Args.push_back(*Iter++);
6076 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6077 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6078 NewCall = Builder.CreateCall(NewFn, Args);
6082 case Intrinsic::bitreverse:
6083 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6086 case Intrinsic::ctlz:
6087 case Intrinsic::cttz: {
6094 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6098 case Intrinsic::objectsize: {
6099 Value *NullIsUnknownSize =
6103 NewCall = Builder.CreateCall(
6108 case Intrinsic::ctpop:
6109 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6111 case Intrinsic::dbg_value: {
6113 Name = Name.substr(5);
6115 if (Name.starts_with(
"dbg.addr")) {
6129 if (
Offset->isNullValue()) {
6130 NewCall = Builder.CreateCall(
6139 case Intrinsic::ptr_annotation:
6147 NewCall = Builder.CreateCall(
6156 case Intrinsic::var_annotation:
6163 NewCall = Builder.CreateCall(
6172 case Intrinsic::riscv_aes32dsi:
6173 case Intrinsic::riscv_aes32dsmi:
6174 case Intrinsic::riscv_aes32esi:
6175 case Intrinsic::riscv_aes32esmi:
6176 case Intrinsic::riscv_sm4ks:
6177 case Intrinsic::riscv_sm4ed: {
6187 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6188 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6194 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6195 Value *Res = NewCall;
6197 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6203 case Intrinsic::nvvm_mapa_shared_cluster: {
6207 Value *Res = NewCall;
6208 Res = Builder.CreateAddrSpaceCast(
6215 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6217 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6219 Args[0] = Builder.CreateAddrSpaceCast(
6223 Args.push_back(Builder.getInt32(0));
6225 NewCall = Builder.CreateCall(NewFn, Args);
6231 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6236 for (
unsigned I = 0;
I < 4; ++
I)
6238 Args.push_back(Builder.getInt32(0));
6239 Args.push_back(Builder.getInt32(0));
6242 Args.push_back(Builder.getInt1(
false));
6243 Args.push_back(Builder.getInt32(0));
6245 NewCall = Builder.CreateCall(NewFn, Args);
6251 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6254 Args[0] = Builder.CreateAddrSpaceCast(
6257 NewCall = Builder.CreateCall(NewFn, Args);
6264#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6265 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6267#undef G2S_CLUSTER_CASE
6272 Args[0] = Builder.CreateAddrSpaceCast(
6278 Args.push_back(Builder.getInt32(0));
6280 NewCall = Builder.CreateCall(NewFn, Args);
6287#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6288 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6296 "expected only the trailing flag_valid_pattern to be missing");
6297 Args.push_back(Builder.getInt32(0));
6299 NewCall = Builder.CreateCall(NewFn, Args);
6305#undef NVVM_TMA_G2S_MODES
6308 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6309 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6310 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6311 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6312 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6313 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6314 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6315 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6317 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6321 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6322 NewCall = Builder.CreateCall(NewFn, Args);
6325 case Intrinsic::nvvm_mbarrier_init: {
6329 if (Args.size() == 2)
6330 Args.push_back(Builder.getInt32(0));
6331 NewCall = Builder.CreateCall(NewFn, Args);
6334 case Intrinsic::riscv_sha256sig0:
6335 case Intrinsic::riscv_sha256sig1:
6336 case Intrinsic::riscv_sha256sum0:
6337 case Intrinsic::riscv_sha256sum1:
6338 case Intrinsic::riscv_sm3p0:
6339 case Intrinsic::riscv_sm3p1: {
6346 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6348 NewCall = Builder.CreateCall(NewFn, Arg);
6350 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6357 case Intrinsic::x86_xop_vfrcz_ss:
6358 case Intrinsic::x86_xop_vfrcz_sd:
6359 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6362 case Intrinsic::x86_xop_vpermil2pd:
6363 case Intrinsic::x86_xop_vpermil2ps:
6364 case Intrinsic::x86_xop_vpermil2pd_256:
6365 case Intrinsic::x86_xop_vpermil2ps_256: {
6369 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6370 NewCall = Builder.CreateCall(NewFn, Args);
6374 case Intrinsic::x86_sse41_ptestc:
6375 case Intrinsic::x86_sse41_ptestz:
6376 case Intrinsic::x86_sse41_ptestnzc: {
6390 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6391 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6393 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6397 case Intrinsic::x86_rdtscp: {
6403 NewCall = Builder.CreateCall(NewFn);
6405 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6408 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6416 case Intrinsic::x86_sse41_insertps:
6417 case Intrinsic::x86_sse41_dppd:
6418 case Intrinsic::x86_sse41_dpps:
6419 case Intrinsic::x86_sse41_mpsadbw:
6420 case Intrinsic::x86_avx_dp_ps_256:
6421 case Intrinsic::x86_avx2_mpsadbw: {
6427 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6428 NewCall = Builder.CreateCall(NewFn, Args);
6432 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6433 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6434 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6435 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6436 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6437 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6443 NewCall = Builder.CreateCall(NewFn, Args);
6452 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6453 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6454 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6455 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6456 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6457 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6461 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6462 Args[1] = Builder.CreateBitCast(
6465 NewCall = Builder.CreateCall(NewFn, Args);
6466 Value *Res = Builder.CreateBitCast(
6474 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6475 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6476 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6480 Args[1] = Builder.CreateBitCast(
6482 Args[2] = Builder.CreateBitCast(
6485 NewCall = Builder.CreateCall(NewFn, Args);
6489 case Intrinsic::amdgcn_asyncmark:
6490 case Intrinsic::amdgcn_wait_asyncmark: {
6494 Args.push_back(Builder.getInt32(0));
6499 NewCall = Builder.CreateCall(NewFn, Args, Bundles);
6507 case Intrinsic::thread_pointer: {
6508 NewCall = Builder.CreateCall(NewFn, {});
6512 case Intrinsic::memcpy:
6513 case Intrinsic::memmove:
6514 case Intrinsic::memset: {
6530 NewCall = Builder.CreateCall(NewFn, Args);
6533 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6534 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6535 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6540 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6543 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6547 case Intrinsic::masked_load:
6548 case Intrinsic::masked_gather:
6549 case Intrinsic::masked_store:
6550 case Intrinsic::masked_scatter: {
6556 auto GetMaybeAlign = [](
Value *
Op) {
6558 uint64_t Val = CI->getZExtValue();
6566 auto GetAlign = [&](
Value *
Op) {
6575 case Intrinsic::masked_load:
6576 NewCall = Builder.CreateMaskedLoad(
6580 case Intrinsic::masked_gather:
6581 NewCall = Builder.CreateMaskedGather(
6587 case Intrinsic::masked_store:
6588 NewCall = Builder.CreateMaskedStore(
6592 case Intrinsic::masked_scatter:
6593 NewCall = Builder.CreateMaskedScatter(
6595 DL.getValueOrABITypeAlignment(
6609 case Intrinsic::lifetime_start:
6610 case Intrinsic::lifetime_end: {
6622 NewCall = Builder.CreateLifetimeStart(Ptr);
6624 NewCall = Builder.CreateLifetimeEnd(Ptr);
6633 case Intrinsic::x86_avx512_vpdpbusd_128:
6634 case Intrinsic::x86_avx512_vpdpbusd_256:
6635 case Intrinsic::x86_avx512_vpdpbusd_512:
6636 case Intrinsic::x86_avx512_vpdpbusds_128:
6637 case Intrinsic::x86_avx512_vpdpbusds_256:
6638 case Intrinsic::x86_avx512_vpdpbusds_512:
6639 case Intrinsic::x86_avx2_vpdpbssd_128:
6640 case Intrinsic::x86_avx2_vpdpbssd_256:
6641 case Intrinsic::x86_avx10_vpdpbssd_512:
6642 case Intrinsic::x86_avx2_vpdpbssds_128:
6643 case Intrinsic::x86_avx2_vpdpbssds_256:
6644 case Intrinsic::x86_avx10_vpdpbssds_512:
6645 case Intrinsic::x86_avx2_vpdpbsud_128:
6646 case Intrinsic::x86_avx2_vpdpbsud_256:
6647 case Intrinsic::x86_avx10_vpdpbsud_512:
6648 case Intrinsic::x86_avx2_vpdpbsuds_128:
6649 case Intrinsic::x86_avx2_vpdpbsuds_256:
6650 case Intrinsic::x86_avx10_vpdpbsuds_512:
6651 case Intrinsic::x86_avx2_vpdpbuud_128:
6652 case Intrinsic::x86_avx2_vpdpbuud_256:
6653 case Intrinsic::x86_avx10_vpdpbuud_512:
6654 case Intrinsic::x86_avx2_vpdpbuuds_128:
6655 case Intrinsic::x86_avx2_vpdpbuuds_256:
6656 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6661 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6662 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6664 NewCall = Builder.CreateCall(NewFn, Args);
6667 case Intrinsic::x86_avx512_vpdpwssd_128:
6668 case Intrinsic::x86_avx512_vpdpwssd_256:
6669 case Intrinsic::x86_avx512_vpdpwssd_512:
6670 case Intrinsic::x86_avx512_vpdpwssds_128:
6671 case Intrinsic::x86_avx512_vpdpwssds_256:
6672 case Intrinsic::x86_avx512_vpdpwssds_512:
6673 case Intrinsic::x86_avx2_vpdpwsud_128:
6674 case Intrinsic::x86_avx2_vpdpwsud_256:
6675 case Intrinsic::x86_avx10_vpdpwsud_512:
6676 case Intrinsic::x86_avx2_vpdpwsuds_128:
6677 case Intrinsic::x86_avx2_vpdpwsuds_256:
6678 case Intrinsic::x86_avx10_vpdpwsuds_512:
6679 case Intrinsic::x86_avx2_vpdpwusd_128:
6680 case Intrinsic::x86_avx2_vpdpwusd_256:
6681 case Intrinsic::x86_avx10_vpdpwusd_512:
6682 case Intrinsic::x86_avx2_vpdpwusds_128:
6683 case Intrinsic::x86_avx2_vpdpwusds_256:
6684 case Intrinsic::x86_avx10_vpdpwusds_512:
6685 case Intrinsic::x86_avx2_vpdpwuud_128:
6686 case Intrinsic::x86_avx2_vpdpwuud_256:
6687 case Intrinsic::x86_avx10_vpdpwuud_512:
6688 case Intrinsic::x86_avx2_vpdpwuuds_128:
6689 case Intrinsic::x86_avx2_vpdpwuuds_256:
6690 case Intrinsic::x86_avx10_vpdpwuuds_512:
6695 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6696 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6698 NewCall = Builder.CreateCall(NewFn, Args);
6701 assert(NewCall &&
"Should have either set this variable or returned through "
6702 "the default case");
6709 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6723 F->eraseFromParent();
6729 if (NumOperands == 0)
6737 if (NumOperands == 3) {
6741 Metadata *Elts2[] = {ScalarType, ScalarType,
6757 if (NumOperands == 0 || NumOperands % 3 != 0)
6762 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6767 if (Upgraded ==
Tag)
6777 if (
Opc != Instruction::BitCast)
6781 Type *SrcTy = V->getType();
6798 if (
Opc != Instruction::BitCast)
6801 Type *SrcTy =
C->getType();
6818 if (Flag.getNumOperands() < 3)
6819 return std::nullopt;
6821 return Name->getString();
6822 return std::nullopt;
6828 if (!
getIROptions(M.getContext()).auto_upgrade_debug_info)
6836 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6837 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6838 if (auto Name = getModuleFlagNameSafely(*Flag))
6839 return *Name ==
"Debug Info Version";
6842 if (OpIt != ModFlags->op_end()) {
6843 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6850 bool BrokenDebugInfo =
false;
6853 if (!BrokenDebugInfo)
6859 M.getContext().diagnose(Diag);
6866 M.getContext().diagnose(DiagVersion);
6876 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6879 if (
F->hasFnAttribute(Attr)) {
6882 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6884 auto [Part, Rest] = S.
split(
',');
6890 const unsigned Dim = DimC -
'x';
6891 assert(Dim < 3 &&
"Unexpected dim char");
6901 F->addFnAttr(Attr, NewAttr);
6905 return S ==
"x" || S ==
"y" || S ==
"z";
6910 if (
K ==
"kernel") {
6922 const unsigned Idx = (AlignIdxValuePair >> 16);
6923 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6928 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6933 if (
K ==
"minctasm") {
6938 if (
K ==
"maxnreg") {
6943 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6947 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6951 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6955 if (
K ==
"grid_constant") {
6970 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6977 if (!SeenNodes.
insert(MD).second)
6984 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6991 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6993 const MDOperand &V = MD->getOperand(j + 1);
6999 if (NewOperands.
size() > 1)
7012 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
7013 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
7014 if (ModRetainReleaseMarker) {
7020 ID->getString().split(ValueComp,
"#");
7021 if (ValueComp.
size() == 2) {
7022 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
7026 M.eraseNamedMetadata(ModRetainReleaseMarker);
7037 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
7063 bool InvalidCast =
false;
7065 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7078 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7080 Args.push_back(Arg);
7087 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7092 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7105 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7113 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7114 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7115 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7116 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7117 {
"objc_autoreleaseReturnValue",
7118 llvm::Intrinsic::objc_autoreleaseReturnValue},
7119 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7120 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7121 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7122 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7123 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7124 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7125 {
"objc_release", llvm::Intrinsic::objc_release},
7126 {
"objc_retain", llvm::Intrinsic::objc_retain},
7127 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7128 {
"objc_retainAutoreleaseReturnValue",
7129 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7130 {
"objc_retainAutoreleasedReturnValue",
7131 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7132 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7133 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7134 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7135 {
"objc_unsafeClaimAutoreleasedReturnValue",
7136 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7137 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7138 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7139 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7140 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7141 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7142 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7143 {
"objc_arc_annotation_topdown_bbstart",
7144 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7145 {
"objc_arc_annotation_topdown_bbend",
7146 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7147 {
"objc_arc_annotation_bottomup_bbstart",
7148 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7149 {
"objc_arc_annotation_bottomup_bbend",
7150 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7152 for (
auto &
I : RuntimeFuncs)
7153 UpgradeToIntrinsic(
I.first,
I.second);
7177 std::optional<bool> UseAddressDisc;
7180 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7181 for (
const MDNode *Flag : ModFlags->operands()) {
7183 if (Name && (*Name ==
"ptrauth-init-fini" ||
7184 *Name ==
"ptrauth-init-fini-address-discrimination"))
7189 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7190 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7191 constexpr unsigned ExpectedAddressMarker = 1;
7194 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7197 bool HasAddressDisc;
7198 if (!CPA->hasAddressDiscriminator())
7199 HasAddressDisc =
false;
7200 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7201 HasAddressDisc =
true;
7205 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7208 UseAddressDisc = HasAddressDisc;
7209 return CPA->getPointer();
7213 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7216 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7218 if (!GV || !GV->hasInitializer())
7222 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7225 std::vector<Constant *> NewStructors;
7226 NewStructors.reserve(OldStructorsArray->getNumOperands());
7228 for (
Use &U : OldStructorsArray->operands()) {
7237 Func = UpgradeSinglePointer(Func);
7241 NewStructors.push_back(
7250 if (GlobalArraysToUpgrade.
empty())
7252 assert(UseAddressDisc.has_value());
7254 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7255 GV->setInitializer(NewInit);
7258 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7268 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7272 bool HasObjCFlag =
false, HasClassProperties =
false;
7273 bool HasSwiftVersionFlag =
false;
7274 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7281 if (
Op->getNumOperands() != 3)
7295 if (ID->getString() ==
"Objective-C Image Info Version")
7297 if (ID->getString() ==
"Objective-C Class Properties")
7298 HasClassProperties =
true;
7300 if (ID->getString() ==
"PIC Level") {
7301 if (
auto *Behavior =
7303 uint64_t V = Behavior->getLimitedValue();
7309 if (ID->getString() ==
"PIE Level")
7310 if (
auto *Behavior =
7318 if (ID->getString() ==
"branch-target-enforcement" ||
7319 (ID->getString().starts_with(
"sign-return-address") &&
7320 ID->getString() !=
"sign-return-address-harden")) {
7321 if (
auto *Behavior =
7327 Op->getOperand(1),
Op->getOperand(2)};
7337 if (ID->getString() ==
"Objective-C Image Info Section") {
7340 Value->getString().split(ValueComp,
" ");
7341 if (ValueComp.
size() != 1) {
7342 std::string NewValue;
7343 for (
auto &S : ValueComp)
7344 NewValue += S.str();
7355 if (ID->getString() ==
"Objective-C Garbage Collection") {
7358 assert(Md->getValue() &&
"Expected non-empty metadata");
7359 auto Type = Md->getValue()->getType();
7362 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7363 if ((Val & 0xff) != Val) {
7364 HasSwiftVersionFlag =
true;
7365 SwiftABIVersion = (Val & 0xff00) >> 8;
7366 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7367 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7378 if (ID->getString() ==
"amdgpu_code_object_version") {
7381 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7390 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7419 if (HasObjCFlag && !HasClassProperties) {
7425 if (HasSwiftVersionFlag) {
7429 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7431 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7439 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7443 auto MatchesVersion = [](
const MDNode *
Op) {
7444 return Op->getNumOperands() >= 3 &&
7458 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7459 assert(
Op->getNumOperands() >= 2 &&
7460 "Expected at least 2 operands - name and linkage type");
7472 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7483 auto TrimSpaces = [](
StringRef Section) -> std::string {
7485 Section.split(Components,
',');
7490 for (
auto Component : Components)
7491 OS <<
',' << Component.trim();
7496 for (
auto &GV : M.globals()) {
7497 if (!GV.hasSection())
7502 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7507 GV.setSection(TrimSpaces(Section));
7523struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7524 StrictFPUpgradeVisitor() =
default;
7527 if (!
Call.isStrictFP())
7533 Call.removeFnAttr(Attribute::StrictFP);
7534 Call.addFnAttr(Attribute::NoBuiltin);
7539struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7540 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7541 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7543 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7558 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7559 StrictFPUpgradeVisitor SFPV;
7564 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7565 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7566 for (
auto &Arg :
F.args())
7568 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7570 bool AddingAttrs =
false, RemovingAttrs =
false;
7571 AttrBuilder AttrsToAdd(
F.getContext());
7576 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7577 A.isValid() &&
A.isStringAttribute()) {
7578 F.setSection(
A.getValueAsString());
7580 RemovingAttrs =
true;
7584 A.isValid() &&
A.isStringAttribute()) {
7587 AddingAttrs = RemovingAttrs =
true;
7590 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7591 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7593 RemovingAttrs =
true;
7594 if (
A.getValueAsString() ==
"true") {
7595 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7604 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7607 if (
A.getValueAsBool()) {
7608 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7614 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7615 RemovingAttrs =
true;
7622 bool HandleDenormalMode =
false;
7624 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7627 DenormalFPMath = ParsedMode;
7629 AddingAttrs = RemovingAttrs =
true;
7630 HandleDenormalMode =
true;
7634 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7638 DenormalFPMathF32 = ParsedMode;
7640 AddingAttrs = RemovingAttrs =
true;
7641 HandleDenormalMode =
true;
7645 if (HandleDenormalMode)
7646 AttrsToAdd.addDenormalFPEnvAttr(
7650 F.removeFnAttrs(AttrsToRemove);
7653 F.addFnAttrs(AttrsToAdd);
7659 if (!
F.hasFnAttribute(FnAttrName)) {
7660 F.addFnAttr(FnAttrName,
Value);
7662 <<
"\", function: " <<
F.getName() <<
"\n");
7670 if (!
F.hasFnAttribute(FnAttrName)) {
7672 F.addFnAttr(FnAttrName);
7674 <<
", function: " <<
F.getName() <<
"\n");
7677 auto A =
F.getFnAttribute(FnAttrName);
7678 if (
"false" ==
A.getValueAsString()) {
7679 F.removeFnAttr(FnAttrName);
7681 <<
"=\"false\", function: " <<
F.getName() <<
"\n");
7682 }
else if (
"true" ==
A.getValueAsString()) {
7683 F.removeFnAttr(FnAttrName);
7684 F.addFnAttr(FnAttrName);
7686 <<
"=\"true\", function: " <<
F.getName() <<
"\n");
7693 M.setModuleFlag(Behavior,
Key, Val);
7694 LLVM_DEBUG(
dbgs() <<
"Converted module flag: " <<
"{" << Behavior <<
", "
7695 <<
Key <<
", " << Val <<
"}\n");
7699 Triple T(M.getTargetTriple());
7700 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7703 uint64_t BTEValue = 0;
7704 uint64_t BPPLRValue = 0;
7705 uint64_t GCSValue = 0;
7706 uint64_t SRAValue = 0;
7707 uint64_t SRAALLValue = 0;
7708 uint64_t SRABKeyValue = 0;
7710 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7714 if (
Op->getNumOperands() != 3)
7723 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7724 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7725 : IDStr ==
"guarded-control-stack" ? &GCSValue
7726 : IDStr ==
"sign-return-address" ? &SRAValue
7727 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7728 : IDStr ==
"sign-return-address-with-bkey"
7734 *ValPtr = CI->getZExtValue();
7738 LLVM_DEBUG(
dbgs() <<
"Found module flag: " << IDStr <<
"(" << *ValPtr
7743 bool BTE = BTEValue == 1;
7744 bool BPPLR = BPPLRValue == 1;
7745 bool GCS = GCSValue == 1;
7746 bool SRA = SRAValue == 1;
7749 if (SRA && SRAALLValue == 1)
7750 SignTypeValue =
"all";
7753 if (SRA && SRABKeyValue == 1)
7754 SignKeyValue =
"b_key";
7756 for (
Function &
F : M.getFunctionList()) {
7757 if (
F.isDeclaration())
7764 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7765 A.isValid() &&
"none" ==
A.getValueAsString()) {
7766 F.removeFnAttr(
"sign-return-address");
7767 F.removeFnAttr(
"sign-return-address-key");
7783 if (SRAALLValue == 1)
7785 if (SRABKeyValue == 1) {
7814 if (
T->getNumOperands() < 1)
7819 if (S->getString().starts_with(
"llvm.vectorizer."))
7825 StringRef OldPrefix =
"llvm.vectorizer.";
7828 if (OldTag ==
"llvm.vectorizer.unroll")
7840 if (
T->getNumOperands() < 1)
7852 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7865 Ops.reserve(
T->getNumOperands());
7866 Ops.push_back(NewTag);
7867 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7868 Ops.push_back(
T->getOperand(
I));
7885 if (
T->isDistinct()) {
7886 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7898 Ops.reserve(
T->getNumOperands());
7909 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7910 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7911 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7914 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7916 auto I =
DL.find(
"-n64-");
7918 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7923 std::string Res =
DL.str();
7926 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7927 Res.append(Res.empty() ?
"G1" :
"-G1");
7935 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7936 Res.append(
"-ni:7:8:9");
7938 if (
DL.ends_with(
"ni:7"))
7940 if (
DL.ends_with(
"ni:7:8"))
7945 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7946 Res.append(
"-p7:160:256:256:32");
7947 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7948 Res.append(
"-p8:128:128:128:48");
7949 constexpr StringRef OldP8(
"-p8:128:128-");
7950 if (
DL.contains(OldP8))
7951 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7952 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7953 Res.append(
"-p9:192:256:256:32");
7958 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7959 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7960 Res.append((
"-" + AS +
":32:32").str());
7965 if (!
DL.contains(
"m:e"))
7966 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7971 if (
T.isSystemZ() && !
DL.empty()) {
7973 if (!
DL.contains(
"-S64"))
7974 return "E-S64" +
DL.drop_front(1).str();
7978 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7981 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7982 if (!
DL.contains(AddrSpaces)) {
7984 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7985 if (R.match(Res, &
Groups))
7991 if (
T.isAArch64()) {
7993 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7994 Res.append(
"-Fn32");
7995 AddPtr32Ptr64AddrSpaces();
7999 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
8003 std::string I64 =
"-i64:64";
8004 std::string I128 =
"-i128:128";
8006 size_t Pos = Res.find(I64);
8007 if (Pos !=
size_t(-1))
8008 Res.insert(Pos + I64.size(), I128);
8012 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
8013 size_t Pos = Res.find(
"-S128");
8016 Res.insert(Pos,
"-f64:32:64");
8021 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
8023 size_t Pos = Res.
find(p3232);
8025 Res.insert(Pos + p3232.
size(),
"-Fi8");
8031 AddPtr32Ptr64AddrSpaces();
8039 if (!
T.isOSIAMCU()) {
8040 std::string I128 =
"-i128:128";
8043 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
8044 if (R.match(Res, &
Groups))
8052 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
8054 auto I =
Ref.find(
"-f80:32-");
8056 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
8064 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
8067 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
8068 B.removeAttribute(
"no-frame-pointer-elim");
8070 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
8072 if (FramePointer !=
"all")
8073 FramePointer =
"non-leaf";
8074 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
8076 if (!FramePointer.
empty())
8077 B.addAttribute(
"frame-pointer", FramePointer);
8079 A =
B.getAttribute(
"null-pointer-is-valid");
8082 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
8083 B.removeAttribute(
"null-pointer-is-valid");
8084 if (NullPointerIsValid)
8085 B.addAttribute(Attribute::NullPointerIsValid);
8088 A =
B.getAttribute(
"uniform-work-group-size");
8092 bool IsTrue = Val ==
"true";
8093 B.removeAttribute(
"uniform-work-group-size");
8095 B.addAttribute(
"uniform-work-group-size");
8106 return OBD.
getTag() ==
"clang.arc.attachedcall" &&
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
AMDGPU address space definition.
AMDGPU Register Bank Select
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
This file contains the simple types necessary to represent the attributes associated with functions a...
static Value * upgradeX86VPERMT2Intrinsics(IRBuilder<> &Builder, CallBase &CI, bool ZeroMask, bool IndexForm)
static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID)
static unsigned getFullArgCountForDefaultArgUpgrade(Function *F, Intrinsic::ID IID, SmallVectorImpl< Type * > &OverloadTys)
#define G2S_ID(ID_SUFFIX, NAME)
static Metadata * upgradeLoopArgument(Metadata *MD)
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name)
static bool isXYZ(StringRef S)
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords)
static Value * upgradeX86PSLLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F, StringRef Name)
static Value * upgradeVPIntrinsicCall(StringRef Name, CallBase *CI, IRBuilder<> &Builder)
static std::optional< unsigned > getNVPTXTMAReductionOp(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name)
static bool upgradeRetainReleaseMarker(Module &M)
This checks for objc retain release marker which should be upgraded.
static Value * upgradeX86vpcom(IRBuilder<> &Builder, CallBase &CI, unsigned Imm, bool IsSigned)
static Value * upgradeMaskToInt(IRBuilder<> &Builder, CallBase &CI)
static bool convertIntrinsicValidType(StringRef Name, const FunctionType *FuncTy)
static Value * upgradeX86Rotate(IRBuilder<> &Builder, CallBase &CI, bool IsRotateRight)
static bool upgradeX86MultiplyAddBytes(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Intrinsic::ID getFunctionalIntrinsicIDForVP(StringRef Name)
static void setFunctionAttrIfNotSet(Function &F, StringRef FnAttrName, StringRef Value)
static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name)
static bool upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K, const Metadata *V)
static MDNode * unwrapMAVOp(CallBase *CI, unsigned Op)
Helper to unwrap intrinsic call MetadataAsValue operands.
static MDString * upgradeLoopTag(LLVMContext &C, StringRef OldTag)
static bool isValidLegacyNVVMFPArithIntrinsic(const Intrinsic::ID(&IIDs)[2][3], StringRef Type, bool IsRN, bool IsFTZ, bool IsSat, bool IsRelu)
static ICmpInst::Predicate getVPIntPredicateFromMD(const Value *Op)
static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC, GlobalValue *GV, const Metadata *V)
static bool upgradeX86MaskedFPCompare(Function *F, Intrinsic::ID IID, Function *&NewFn)
static Value * upgradeX86ALIGNIntrinsics(IRBuilder<> &Builder, Value *Op0, Value *Op1, Value *Shift, Value *Passthru, Value *Mask, bool IsVALIGN)
static Value * upgradeAbs(IRBuilder<> &Builder, CallBase &CI)
static bool shouldUpgradeVPIntrinsic(StringRef Name)
static Value * emitX86Select(IRBuilder<> &Builder, Value *Mask, Value *Op0, Value *Op1)
static Value * upgradeAArch64IntrinsicCall(StringRef Name, CallBase *CI, Function *F, IRBuilder<> &Builder)
#define G2S_CTA_ID(ID_SUFFIX, NAME)
static Value * upgradeMaskedMove(IRBuilder<> &Builder, CallBase &CI)
static constexpr Intrinsic::ID NVVMFMAIIDs[2][3]
static constexpr Intrinsic::ID NVVMFMulIIDs[2][3]
static const BooleanLoopTags * getOldBooleanLoopTags(const MDTuple *T)
Return the replacement tags if T still uses a removed two-operand form.
static bool upgradeX86IntrinsicFunction(Function *F, StringRef Name, Function *&NewFn)
static Value * applyX86MaskOn1BitsVec(IRBuilder<> &Builder, Value *Vec, Value *Mask)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static 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 * upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI, StringRef Name, const Intrinsic::ID(&IIDs)[2][3])
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 Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocSharedIntrinsic(Function *F, StringRef Name)
static bool upgradeAVX512MaskToSelect(StringRef Name, IRBuilder<> &Builder, CallBase &CI, Value *&Rep)
static void upgradeDbgIntrinsicToDbgRecord(StringRef Name, CallBase *CI)
Convert debug intrinsic calls to non-instruction debug records.
static void ConvertFunctionAttr(Function &F, bool Set, StringRef FnAttrName)
static Value * upgradePMULDQ(IRBuilder<> &Builder, CallBase &CI, bool IsSigned)
static std::optional< std::pair< Intrinsic::ID, RoundingMode > > getNVVMFPArithUpgrade(StringRef Name, const Intrinsic::ID(&IIDs)[2][3])
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 constexpr Intrinsic::ID NVVMFAddIIDs[2][3]
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 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...
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.