36#include "llvm/IR/IntrinsicsAArch64.h"
37#include "llvm/IR/IntrinsicsAMDGPU.h"
38#include "llvm/IR/IntrinsicsARM.h"
39#include "llvm/IR/IntrinsicsNVPTX.h"
40#include "llvm/IR/IntrinsicsRISCV.h"
41#include "llvm/IR/IntrinsicsWebAssembly.h"
42#include "llvm/IR/IntrinsicsX86.h"
65#define DEBUG_TYPE "auto-upgrade"
84 Type *Arg0Type =
F->getFunctionType()->getParamType(0);
99 Type *LastArgType =
F->getFunctionType()->getParamType(
100 F->getFunctionType()->getNumParams() - 1);
115 if (
F->getReturnType()->isVectorTy())
128 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
129 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
146 Type *Arg1Type =
F->getFunctionType()->getParamType(1);
147 Type *Arg2Type =
F->getFunctionType()->getParamType(2);
161 if (
F->getReturnType()->getScalarType()->isBFloatTy())
171 if (
F->getFunctionType()->getParamType(1)->getScalarType()->isBFloatTy())
185 if (Name.consume_front(
"avx."))
186 return (Name.starts_with(
"blend.p") ||
187 Name ==
"cvt.ps2.pd.256" ||
188 Name ==
"cvtdq2.pd.256" ||
189 Name ==
"cvtdq2.ps.256" ||
190 Name.starts_with(
"movnt.") ||
191 Name.starts_with(
"sqrt.p") ||
192 Name.starts_with(
"storeu.") ||
193 Name.starts_with(
"vbroadcast.s") ||
194 Name.starts_with(
"vbroadcastf128") ||
195 Name.starts_with(
"vextractf128.") ||
196 Name.starts_with(
"vinsertf128.") ||
197 Name.starts_with(
"vperm2f128.") ||
198 Name.starts_with(
"vpermil."));
200 if (Name.consume_front(
"avx2."))
201 return (Name ==
"movntdqa" ||
202 Name.starts_with(
"pabs.") ||
203 Name.starts_with(
"padds.") ||
204 Name.starts_with(
"paddus.") ||
205 Name.starts_with(
"pblendd.") ||
207 Name.starts_with(
"pbroadcast") ||
208 Name.starts_with(
"pcmpeq.") ||
209 Name.starts_with(
"pcmpgt.") ||
210 Name.starts_with(
"pmax") ||
211 Name.starts_with(
"pmin") ||
212 Name.starts_with(
"pmovsx") ||
213 Name.starts_with(
"pmovzx") ||
214 Name.starts_with(
"pmulh.w") ||
215 Name.starts_with(
"pmulhu.w") ||
217 Name ==
"pmulu.dq" ||
218 Name.starts_with(
"psll.dq") ||
219 Name.starts_with(
"psrl.dq") ||
220 Name.starts_with(
"psubs.") ||
221 Name.starts_with(
"psubus.") ||
222 Name.starts_with(
"vbroadcast") ||
223 Name ==
"vbroadcasti128" ||
224 Name ==
"vextracti128" ||
225 Name ==
"vinserti128" ||
226 Name ==
"vperm2i128");
228 if (Name.consume_front(
"avx512.")) {
229 if (Name.consume_front(
"mask."))
231 return (Name.starts_with(
"add.p") ||
232 Name.starts_with(
"and.") ||
233 Name.starts_with(
"andn.") ||
234 Name.starts_with(
"broadcast.s") ||
235 Name.starts_with(
"broadcastf32x4.") ||
236 Name.starts_with(
"broadcastf32x8.") ||
237 Name.starts_with(
"broadcastf64x2.") ||
238 Name.starts_with(
"broadcastf64x4.") ||
239 Name.starts_with(
"broadcasti32x4.") ||
240 Name.starts_with(
"broadcasti32x8.") ||
241 Name.starts_with(
"broadcasti64x2.") ||
242 Name.starts_with(
"broadcasti64x4.") ||
243 Name.starts_with(
"cmp.b") ||
244 Name.starts_with(
"cmp.d") ||
245 Name.starts_with(
"cmp.q") ||
246 Name.starts_with(
"cmp.w") ||
247 Name.starts_with(
"compress.b") ||
248 Name.starts_with(
"compress.d") ||
249 Name.starts_with(
"compress.p") ||
250 Name.starts_with(
"compress.q") ||
251 Name.starts_with(
"compress.store.") ||
252 Name.starts_with(
"compress.w") ||
253 Name.starts_with(
"conflict.") ||
254 Name.starts_with(
"cvtdq2pd.") ||
255 Name.starts_with(
"cvtdq2ps.") ||
256 Name ==
"cvtpd2dq.256" ||
257 Name ==
"cvtpd2ps.256" ||
258 Name ==
"cvtps2pd.128" ||
259 Name ==
"cvtps2pd.256" ||
260 Name.starts_with(
"cvtqq2pd.") ||
261 Name ==
"cvtqq2ps.256" ||
262 Name ==
"cvtqq2ps.512" ||
263 Name ==
"cvttpd2dq.256" ||
264 Name ==
"cvttps2dq.128" ||
265 Name ==
"cvttps2dq.256" ||
266 Name.starts_with(
"cvtudq2pd.") ||
267 Name.starts_with(
"cvtudq2ps.") ||
268 Name.starts_with(
"cvtuqq2pd.") ||
269 Name ==
"cvtuqq2ps.256" ||
270 Name ==
"cvtuqq2ps.512" ||
271 Name.starts_with(
"dbpsadbw.") ||
272 Name.starts_with(
"div.p") ||
273 Name.starts_with(
"expand.b") ||
274 Name.starts_with(
"expand.d") ||
275 Name.starts_with(
"expand.load.") ||
276 Name.starts_with(
"expand.p") ||
277 Name.starts_with(
"expand.q") ||
278 Name.starts_with(
"expand.w") ||
279 Name.starts_with(
"fpclass.p") ||
280 Name.starts_with(
"insert") ||
281 Name.starts_with(
"load.") ||
282 Name.starts_with(
"loadu.") ||
283 Name.starts_with(
"lzcnt.") ||
284 Name.starts_with(
"max.p") ||
285 Name.starts_with(
"min.p") ||
286 Name.starts_with(
"movddup") ||
287 Name.starts_with(
"move.s") ||
288 Name.starts_with(
"movshdup") ||
289 Name.starts_with(
"movsldup") ||
290 Name.starts_with(
"mul.p") ||
291 Name.starts_with(
"or.") ||
292 Name.starts_with(
"pabs.") ||
293 Name.starts_with(
"packssdw.") ||
294 Name.starts_with(
"packsswb.") ||
295 Name.starts_with(
"packusdw.") ||
296 Name.starts_with(
"packuswb.") ||
297 Name.starts_with(
"padd.") ||
298 Name.starts_with(
"padds.") ||
299 Name.starts_with(
"paddus.") ||
300 Name.starts_with(
"palignr.") ||
301 Name.starts_with(
"pand.") ||
302 Name.starts_with(
"pandn.") ||
303 Name.starts_with(
"pavg") ||
304 Name.starts_with(
"pbroadcast") ||
305 Name.starts_with(
"pcmpeq.") ||
306 Name.starts_with(
"pcmpgt.") ||
307 Name.starts_with(
"perm.df.") ||
308 Name.starts_with(
"perm.di.") ||
309 Name.starts_with(
"permvar.") ||
310 Name.starts_with(
"pmaddubs.w.") ||
311 Name.starts_with(
"pmaddw.d.") ||
312 Name.starts_with(
"pmax") ||
313 Name.starts_with(
"pmin") ||
314 Name ==
"pmov.qd.256" ||
315 Name ==
"pmov.qd.512" ||
316 Name ==
"pmov.wb.256" ||
317 Name ==
"pmov.wb.512" ||
318 Name.starts_with(
"pmovsx") ||
319 Name.starts_with(
"pmovzx") ||
320 Name.starts_with(
"pmul.dq.") ||
321 Name.starts_with(
"pmul.hr.sw.") ||
322 Name.starts_with(
"pmulh.w.") ||
323 Name.starts_with(
"pmulhu.w.") ||
324 Name.starts_with(
"pmull.") ||
325 Name.starts_with(
"pmultishift.qb.") ||
326 Name.starts_with(
"pmulu.dq.") ||
327 Name.starts_with(
"por.") ||
328 Name.starts_with(
"prol.") ||
329 Name.starts_with(
"prolv.") ||
330 Name.starts_with(
"pror.") ||
331 Name.starts_with(
"prorv.") ||
332 Name.starts_with(
"pshuf.b.") ||
333 Name.starts_with(
"pshuf.d.") ||
334 Name.starts_with(
"pshufh.w.") ||
335 Name.starts_with(
"pshufl.w.") ||
336 Name.starts_with(
"psll.d") ||
337 Name.starts_with(
"psll.q") ||
338 Name.starts_with(
"psll.w") ||
339 Name.starts_with(
"pslli") ||
340 Name.starts_with(
"psllv") ||
341 Name.starts_with(
"psra.d") ||
342 Name.starts_with(
"psra.q") ||
343 Name.starts_with(
"psra.w") ||
344 Name.starts_with(
"psrai") ||
345 Name.starts_with(
"psrav") ||
346 Name.starts_with(
"psrl.d") ||
347 Name.starts_with(
"psrl.q") ||
348 Name.starts_with(
"psrl.w") ||
349 Name.starts_with(
"psrli") ||
350 Name.starts_with(
"psrlv") ||
351 Name.starts_with(
"psub.") ||
352 Name.starts_with(
"psubs.") ||
353 Name.starts_with(
"psubus.") ||
354 Name.starts_with(
"pternlog.") ||
355 Name.starts_with(
"punpckh") ||
356 Name.starts_with(
"punpckl") ||
357 Name.starts_with(
"pxor.") ||
358 Name.starts_with(
"shuf.f") ||
359 Name.starts_with(
"shuf.i") ||
360 Name.starts_with(
"shuf.p") ||
361 Name.starts_with(
"sqrt.p") ||
362 Name.starts_with(
"store.b.") ||
363 Name.starts_with(
"store.d.") ||
364 Name.starts_with(
"store.p") ||
365 Name.starts_with(
"store.q.") ||
366 Name.starts_with(
"store.w.") ||
367 Name ==
"store.ss" ||
368 Name.starts_with(
"storeu.") ||
369 Name.starts_with(
"sub.p") ||
370 Name.starts_with(
"ucmp.") ||
371 Name.starts_with(
"unpckh.") ||
372 Name.starts_with(
"unpckl.") ||
373 Name.starts_with(
"valign.") ||
374 Name ==
"vcvtph2ps.128" ||
375 Name ==
"vcvtph2ps.256" ||
376 Name.starts_with(
"vextract") ||
377 Name.starts_with(
"vfmadd.") ||
378 Name.starts_with(
"vfmaddsub.") ||
379 Name.starts_with(
"vfnmadd.") ||
380 Name.starts_with(
"vfnmsub.") ||
381 Name.starts_with(
"vpdpbusd.") ||
382 Name.starts_with(
"vpdpbusds.") ||
383 Name.starts_with(
"vpdpwssd.") ||
384 Name.starts_with(
"vpdpwssds.") ||
385 Name.starts_with(
"vpermi2var.") ||
386 Name.starts_with(
"vpermil.p") ||
387 Name.starts_with(
"vpermilvar.") ||
388 Name.starts_with(
"vpermt2var.") ||
389 Name.starts_with(
"vpmadd52") ||
390 Name.starts_with(
"vpshld.") ||
391 Name.starts_with(
"vpshldv.") ||
392 Name.starts_with(
"vpshrd.") ||
393 Name.starts_with(
"vpshrdv.") ||
394 Name.starts_with(
"vpshufbitqmb.") ||
395 Name.starts_with(
"xor."));
397 if (Name.consume_front(
"mask3."))
399 return (Name.starts_with(
"vfmadd.") ||
400 Name.starts_with(
"vfmaddsub.") ||
401 Name.starts_with(
"vfmsub.") ||
402 Name.starts_with(
"vfmsubadd.") ||
403 Name.starts_with(
"vfnmsub."));
405 if (Name.consume_front(
"maskz."))
407 return (Name.starts_with(
"pternlog.") ||
408 Name.starts_with(
"vfmadd.") ||
409 Name.starts_with(
"vfmaddsub.") ||
410 Name.starts_with(
"vpdpbusd.") ||
411 Name.starts_with(
"vpdpbusds.") ||
412 Name.starts_with(
"vpdpwssd.") ||
413 Name.starts_with(
"vpdpwssds.") ||
414 Name.starts_with(
"vpermt2var.") ||
415 Name.starts_with(
"vpmadd52") ||
416 Name.starts_with(
"vpshldv.") ||
417 Name.starts_with(
"vpshrdv."));
420 return (Name ==
"movntdqa" ||
421 Name ==
"pmul.dq.512" ||
422 Name ==
"pmulu.dq.512" ||
423 Name.starts_with(
"broadcastm") ||
424 Name.starts_with(
"cmp.p") ||
425 Name.starts_with(
"cvtb2mask.") ||
426 Name.starts_with(
"cvtd2mask.") ||
427 Name.starts_with(
"cvtmask2") ||
428 Name.starts_with(
"cvtq2mask.") ||
429 Name ==
"cvtusi2sd" ||
430 Name.starts_with(
"cvtw2mask.") ||
435 Name ==
"kortestc.w" ||
436 Name ==
"kortestz.w" ||
437 Name.starts_with(
"kunpck") ||
440 Name.starts_with(
"padds.") ||
441 Name.starts_with(
"pbroadcast") ||
442 Name.starts_with(
"pmulh.w") ||
443 Name.starts_with(
"pmulhu.w") ||
444 Name.starts_with(
"prol") ||
445 Name.starts_with(
"pror") ||
446 Name.starts_with(
"psll.dq") ||
447 Name.starts_with(
"psrl.dq") ||
448 Name.starts_with(
"psubs.") ||
449 Name.starts_with(
"ptestm") ||
450 Name.starts_with(
"ptestnm") ||
451 Name.starts_with(
"storent.") ||
452 Name.starts_with(
"vbroadcast.s") ||
453 Name.starts_with(
"vpshld.") ||
454 Name.starts_with(
"vpshrd."));
457 if (Name.consume_front(
"fma."))
458 return (Name.starts_with(
"vfmadd.") ||
459 Name.starts_with(
"vfmsub.") ||
460 Name.starts_with(
"vfmsubadd.") ||
461 Name.starts_with(
"vfnmadd.") ||
462 Name.starts_with(
"vfnmsub."));
464 if (Name.consume_front(
"fma4."))
465 return Name.starts_with(
"vfmadd.s");
467 if (Name.consume_front(
"sse."))
468 return (Name ==
"add.ss" ||
469 Name ==
"cvtsi2ss" ||
470 Name ==
"cvtsi642ss" ||
473 Name.starts_with(
"sqrt.p") ||
475 Name.starts_with(
"storeu.") ||
478 if (Name.consume_front(
"sse2."))
479 return (Name ==
"add.sd" ||
480 Name ==
"cvtdq2pd" ||
481 Name ==
"cvtdq2ps" ||
482 Name ==
"cvtps2pd" ||
483 Name ==
"cvtsi2sd" ||
484 Name ==
"cvtsi642sd" ||
485 Name ==
"cvtss2sd" ||
488 Name.starts_with(
"padds.") ||
489 Name.starts_with(
"paddus.") ||
490 Name.starts_with(
"pcmpeq.") ||
491 Name.starts_with(
"pcmpgt.") ||
497 Name ==
"pmulhu.w" ||
498 Name ==
"pmulu.dq" ||
499 Name.starts_with(
"pshuf") ||
500 Name.starts_with(
"psll.dq") ||
501 Name.starts_with(
"psrl.dq") ||
502 Name.starts_with(
"psubs.") ||
503 Name.starts_with(
"psubus.") ||
504 Name.starts_with(
"sqrt.p") ||
506 Name ==
"storel.dq" ||
507 Name.starts_with(
"storeu.") ||
510 if (Name.consume_front(
"sse41."))
511 return (Name.starts_with(
"blendp") ||
512 Name ==
"movntdqa" ||
522 Name.starts_with(
"pmovsx") ||
523 Name.starts_with(
"pmovzx") ||
526 if (Name.consume_front(
"sse42."))
527 return Name ==
"crc32.64.8";
529 if (Name.consume_front(
"sse4a."))
530 return Name.starts_with(
"movnt.");
532 if (Name.consume_front(
"ssse3."))
533 return (Name ==
"pabs.b.128" ||
534 Name ==
"pabs.d.128" ||
535 Name ==
"pabs.w.128");
537 if (Name.consume_front(
"xop."))
538 return (Name ==
"vpcmov" ||
539 Name ==
"vpcmov.256" ||
540 Name.starts_with(
"vpcom") ||
541 Name.starts_with(
"vprot"));
543 if (Name.consume_front(
"bmi."))
544 return (Name.starts_with(
"pdep.") ||
545 Name.starts_with(
"pext."));
547 return (Name ==
"addcarry.u32" ||
548 Name ==
"addcarry.u64" ||
549 Name ==
"addcarryx.u32" ||
550 Name ==
"addcarryx.u64" ||
551 Name ==
"subborrow.u32" ||
552 Name ==
"subborrow.u64" ||
553 Name.starts_with(
"vcvtph2ps."));
559 if (!Name.consume_front(
"x86."))
567 if (Name ==
"rdtscp") {
569 if (
F->getFunctionType()->getNumParams() == 0)
574 Intrinsic::x86_rdtscp);
581 if (Name.consume_front(
"sse41.ptest")) {
583 .
Case(
"c", Intrinsic::x86_sse41_ptestc)
584 .
Case(
"z", Intrinsic::x86_sse41_ptestz)
585 .
Case(
"nzc", Intrinsic::x86_sse41_ptestnzc)
598 .
Case(
"sse41.insertps", Intrinsic::x86_sse41_insertps)
599 .
Case(
"sse41.dppd", Intrinsic::x86_sse41_dppd)
600 .
Case(
"sse41.dpps", Intrinsic::x86_sse41_dpps)
601 .
Case(
"sse41.mpsadbw", Intrinsic::x86_sse41_mpsadbw)
602 .
Case(
"avx.dp.ps.256", Intrinsic::x86_avx_dp_ps_256)
603 .
Case(
"avx2.mpsadbw", Intrinsic::x86_avx2_mpsadbw)
608 if (Name.consume_front(
"avx512.")) {
609 if (Name.consume_front(
"mask.cmp.")) {
612 .
Case(
"pd.128", Intrinsic::x86_avx512_mask_cmp_pd_128)
613 .
Case(
"pd.256", Intrinsic::x86_avx512_mask_cmp_pd_256)
614 .
Case(
"pd.512", Intrinsic::x86_avx512_mask_cmp_pd_512)
615 .
Case(
"ps.128", Intrinsic::x86_avx512_mask_cmp_ps_128)
616 .
Case(
"ps.256", Intrinsic::x86_avx512_mask_cmp_ps_256)
617 .
Case(
"ps.512", Intrinsic::x86_avx512_mask_cmp_ps_512)
621 }
else if (Name.starts_with(
"vpdpbusd.") ||
622 Name.starts_with(
"vpdpbusds.")) {
625 .
Case(
"vpdpbusd.128", Intrinsic::x86_avx512_vpdpbusd_128)
626 .
Case(
"vpdpbusd.256", Intrinsic::x86_avx512_vpdpbusd_256)
627 .
Case(
"vpdpbusd.512", Intrinsic::x86_avx512_vpdpbusd_512)
628 .
Case(
"vpdpbusds.128", Intrinsic::x86_avx512_vpdpbusds_128)
629 .
Case(
"vpdpbusds.256", Intrinsic::x86_avx512_vpdpbusds_256)
630 .
Case(
"vpdpbusds.512", Intrinsic::x86_avx512_vpdpbusds_512)
634 }
else if (Name.starts_with(
"vpdpwssd.") ||
635 Name.starts_with(
"vpdpwssds.")) {
638 .
Case(
"vpdpwssd.128", Intrinsic::x86_avx512_vpdpwssd_128)
639 .
Case(
"vpdpwssd.256", Intrinsic::x86_avx512_vpdpwssd_256)
640 .
Case(
"vpdpwssd.512", Intrinsic::x86_avx512_vpdpwssd_512)
641 .
Case(
"vpdpwssds.128", Intrinsic::x86_avx512_vpdpwssds_128)
642 .
Case(
"vpdpwssds.256", Intrinsic::x86_avx512_vpdpwssds_256)
643 .
Case(
"vpdpwssds.512", Intrinsic::x86_avx512_vpdpwssds_512)
651 if (Name.consume_front(
"avx2.")) {
652 if (Name.consume_front(
"vpdpb")) {
655 .
Case(
"ssd.128", Intrinsic::x86_avx2_vpdpbssd_128)
656 .
Case(
"ssd.256", Intrinsic::x86_avx2_vpdpbssd_256)
657 .
Case(
"ssds.128", Intrinsic::x86_avx2_vpdpbssds_128)
658 .
Case(
"ssds.256", Intrinsic::x86_avx2_vpdpbssds_256)
659 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpbsud_128)
660 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpbsud_256)
661 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpbsuds_128)
662 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpbsuds_256)
663 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpbuud_128)
664 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpbuud_256)
665 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpbuuds_128)
666 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpbuuds_256)
670 }
else if (Name.consume_front(
"vpdpw")) {
673 .
Case(
"sud.128", Intrinsic::x86_avx2_vpdpwsud_128)
674 .
Case(
"sud.256", Intrinsic::x86_avx2_vpdpwsud_256)
675 .
Case(
"suds.128", Intrinsic::x86_avx2_vpdpwsuds_128)
676 .
Case(
"suds.256", Intrinsic::x86_avx2_vpdpwsuds_256)
677 .
Case(
"usd.128", Intrinsic::x86_avx2_vpdpwusd_128)
678 .
Case(
"usd.256", Intrinsic::x86_avx2_vpdpwusd_256)
679 .
Case(
"usds.128", Intrinsic::x86_avx2_vpdpwusds_128)
680 .
Case(
"usds.256", Intrinsic::x86_avx2_vpdpwusds_256)
681 .
Case(
"uud.128", Intrinsic::x86_avx2_vpdpwuud_128)
682 .
Case(
"uud.256", Intrinsic::x86_avx2_vpdpwuud_256)
683 .
Case(
"uuds.128", Intrinsic::x86_avx2_vpdpwuuds_128)
684 .
Case(
"uuds.256", Intrinsic::x86_avx2_vpdpwuuds_256)
692 if (Name.consume_front(
"avx10.")) {
693 if (Name.consume_front(
"vpdpb")) {
696 .
Case(
"ssd.512", Intrinsic::x86_avx10_vpdpbssd_512)
697 .
Case(
"ssds.512", Intrinsic::x86_avx10_vpdpbssds_512)
698 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpbsud_512)
699 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpbsuds_512)
700 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpbuud_512)
701 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpbuuds_512)
705 }
else if (Name.consume_front(
"vpdpw")) {
707 .
Case(
"sud.512", Intrinsic::x86_avx10_vpdpwsud_512)
708 .
Case(
"suds.512", Intrinsic::x86_avx10_vpdpwsuds_512)
709 .
Case(
"usd.512", Intrinsic::x86_avx10_vpdpwusd_512)
710 .
Case(
"usds.512", Intrinsic::x86_avx10_vpdpwusds_512)
711 .
Case(
"uud.512", Intrinsic::x86_avx10_vpdpwuud_512)
712 .
Case(
"uuds.512", Intrinsic::x86_avx10_vpdpwuuds_512)
720 if (Name.consume_front(
"avx512bf16.")) {
723 .
Case(
"cvtne2ps2bf16.128",
724 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128)
725 .
Case(
"cvtne2ps2bf16.256",
726 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256)
727 .
Case(
"cvtne2ps2bf16.512",
728 Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512)
729 .
Case(
"mask.cvtneps2bf16.128",
730 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
731 .
Case(
"cvtneps2bf16.256",
732 Intrinsic::x86_avx512bf16_cvtneps2bf16_256)
733 .
Case(
"cvtneps2bf16.512",
734 Intrinsic::x86_avx512bf16_cvtneps2bf16_512)
741 .
Case(
"dpbf16ps.128", Intrinsic::x86_avx512bf16_dpbf16ps_128)
742 .
Case(
"dpbf16ps.256", Intrinsic::x86_avx512bf16_dpbf16ps_256)
743 .
Case(
"dpbf16ps.512", Intrinsic::x86_avx512bf16_dpbf16ps_512)
750 if (Name.consume_front(
"xop.")) {
752 if (Name.starts_with(
"vpermil2")) {
755 auto Idx =
F->getFunctionType()->getParamType(2);
756 if (Idx->isFPOrFPVectorTy()) {
757 unsigned IdxSize = Idx->getPrimitiveSizeInBits();
758 unsigned EltSize = Idx->getScalarSizeInBits();
759 if (EltSize == 64 && IdxSize == 128)
760 ID = Intrinsic::x86_xop_vpermil2pd;
761 else if (EltSize == 32 && IdxSize == 128)
762 ID = Intrinsic::x86_xop_vpermil2ps;
763 else if (EltSize == 64 && IdxSize == 256)
764 ID = Intrinsic::x86_xop_vpermil2pd_256;
766 ID = Intrinsic::x86_xop_vpermil2ps_256;
768 }
else if (
F->arg_size() == 2)
771 .
Case(
"vfrcz.ss", Intrinsic::x86_xop_vfrcz_ss)
772 .
Case(
"vfrcz.sd", Intrinsic::x86_xop_vfrcz_sd)
783 if (Name ==
"seh.recoverfp") {
785 Intrinsic::eh_recoverfp);
797 if (Name.starts_with(
"rbit")) {
800 F->getParent(), Intrinsic::bitreverse,
F->arg_begin()->getType());
804 if (Name ==
"thread.pointer") {
807 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
811 bool Neon = Name.consume_front(
"neon.");
816 if (Name.consume_front(
"bfdot.")) {
820 .
Cases({
"v2f32.v8i8",
"v4f32.v16i8"},
825 size_t OperandWidth =
F->getReturnType()->getPrimitiveSizeInBits();
826 assert((OperandWidth == 64 || OperandWidth == 128) &&
827 "Unexpected operand width");
829 std::array<Type *, 2> Tys{
840 if (Name.consume_front(
"bfm")) {
842 if (Name.consume_back(
".v4f32.v16i8")) {
888 F->arg_begin()->getType());
892 if (Name.consume_front(
"vst")) {
894 static const Regex vstRegex(
"^([1234]|[234]lane)\\.v[a-z0-9]*$");
898 Intrinsic::arm_neon_vst1, Intrinsic::arm_neon_vst2,
899 Intrinsic::arm_neon_vst3, Intrinsic::arm_neon_vst4};
902 Intrinsic::arm_neon_vst2lane, Intrinsic::arm_neon_vst3lane,
903 Intrinsic::arm_neon_vst4lane};
905 auto fArgs =
F->getFunctionType()->params();
906 Type *Tys[] = {fArgs[0], fArgs[1]};
909 F->getParent(), StoreInts[fArgs.size() - 3], Tys);
912 F->getParent(), StoreLaneInts[fArgs.size() - 5], Tys);
921 if (Name.consume_front(
"mve.")) {
923 if (Name ==
"vctp64") {
933 if (Name.starts_with(
"vrintn.v")) {
935 F->getParent(), Intrinsic::roundeven,
F->arg_begin()->getType());
940 if (Name.consume_back(
".v4i1")) {
942 if (Name.consume_back(
".predicated.v2i64.v4i32"))
944 return Name ==
"mull.int" || Name ==
"vqdmull";
946 if (Name.consume_back(
".v2i64")) {
948 bool IsGather = Name.consume_front(
"vldr.gather.");
949 if (IsGather || Name.consume_front(
"vstr.scatter.")) {
950 if (Name.consume_front(
"base.")) {
952 Name.consume_front(
"wb.");
955 return Name ==
"predicated.v2i64";
958 if (Name.consume_front(
"offset.predicated."))
959 return Name == (IsGather ?
"v2i64.p0i64" :
"p0i64.v2i64") ||
960 Name == (IsGather ?
"v2i64.p0" :
"p0.v2i64");
973 if (Name.consume_front(
"cde.vcx")) {
975 if (Name.consume_back(
".predicated.v2i64.v4i1"))
977 return Name ==
"1q" || Name ==
"1qa" || Name ==
"2q" || Name ==
"2qa" ||
978 Name ==
"3q" || Name ==
"3qa";
992 F->arg_begin()->getType());
998 .
Case(
"smax", Intrinsic::smax)
999 .
Case(
"smin", Intrinsic::smin)
1000 .
Case(
"umax", Intrinsic::umax)
1001 .
Case(
"umin", Intrinsic::umin)
1004 if (
F->arg_size() != 2 || !
F->getReturnType()->isIntOrIntVectorTy())
1007 F->getReturnType());
1011 if (Name.starts_with(
"addp")) {
1013 if (
F->arg_size() != 2)
1016 if (Ty && Ty->getElementType()->isFloatingPointTy()) {
1018 F->getParent(), Intrinsic::aarch64_neon_faddp, Ty);
1024 if (Name.starts_with(
"bfcvt")) {
1030 if (Name ==
"vcvtfp2hf" || Name ==
"vcvthf2fp") {
1037 if (Name.consume_front(
"sve.")) {
1039 if (Name.consume_front(
"bf")) {
1040 if (Name ==
"mmla") {
1041 Type *Tys[] = {
F->getReturnType(),
1042 std::next(
F->arg_begin())->getType()};
1044 F->getParent(), Intrinsic::aarch64_sve_fmmla, Tys);
1047 if (Name.consume_back(
".lane")) {
1051 .
Case(
"dot", Intrinsic::aarch64_sve_bfdot_lane_v2)
1052 .
Case(
"mlalb", Intrinsic::aarch64_sve_bfmlalb_lane_v2)
1053 .
Case(
"mlalt", Intrinsic::aarch64_sve_bfmlalt_lane_v2)
1065 if (Name ==
"fcvt.bf16f32" || Name ==
"fcvtnt.bf16f32") {
1070 if (Name.consume_front(
"convert.from.svbool")) {
1073 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1076 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_to_svcount;
1081 if (Name.consume_front(
"convert.to.svbool")) {
1084 if (!TTy || TTy->getName() !=
"aarch64.svcount")
1087 Intrinsic::ID ID = Intrinsic::aarch64_sve_convert_from_svcount;
1092 if (Name.consume_front(
"addqv")) {
1094 if (!
F->getReturnType()->isFPOrFPVectorTy())
1097 auto Args =
F->getFunctionType()->params();
1098 Type *Tys[] = {
F->getReturnType(), Args[1]};
1100 F->getParent(), Intrinsic::aarch64_sve_faddqv, Tys);
1104 if (Name.consume_front(
"ld")) {
1106 static const Regex LdRegex(
"^[234](.nxv[a-z0-9]+|$)");
1107 if (LdRegex.
match(Name)) {
1113 "Expected 2 arguments for ld* intrinsic.");
1114 Type *PtrTy =
F->getArg(1)->getType();
1117 Intrinsic::aarch64_sve_ld2_sret,
1118 Intrinsic::aarch64_sve_ld3_sret,
1119 Intrinsic::aarch64_sve_ld4_sret,
1122 F->getParent(), LoadIDs[Name[0] -
'2'], {Ty, PtrTy});
1128 if (Name.consume_front(
"tuple.")) {
1130 if (Name.starts_with(
"get")) {
1132 Type *Tys[] = {
F->getReturnType(),
F->arg_begin()->getType()};
1134 F->getParent(), Intrinsic::vector_extract, Tys);
1138 if (Name.starts_with(
"set")) {
1140 auto Args =
F->getFunctionType()->params();
1141 Type *Tys[] = {Args[0], Args[2], Args[1]};
1143 F->getParent(), Intrinsic::vector_insert, Tys);
1147 static const Regex CreateTupleRegex(
"^create[234](.nxv[a-z0-9]+|$)");
1148 if (CreateTupleRegex.
match(Name)) {
1150 auto Args =
F->getFunctionType()->params();
1151 Type *Tys[] = {
F->getReturnType(), Args[1]};
1153 F->getParent(), Intrinsic::vector_insert, Tys);
1159 if (Name.starts_with(
"rev.nxv")) {
1162 F->getParent(), Intrinsic::vector_reverse,
F->getReturnType());
1168 if (Name.consume_front(
"sme.")) {
1170 if (Name.consume_front(
"ftmopa.")) {
1175 .
Case(
"za16.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za16)
1176 .
Case(
"za32.nxv16i8", Intrinsic::aarch64_sme_fp8_ftmopa_za32)
1194#define NVVM_TMA_G2S_MODES(M) \
1195 M(tile_1d, "tile.1d") \
1196 M(tile_2d, "tile.2d") \
1197 M(tile_3d, "tile.3d") \
1198 M(tile_4d, "tile.4d") \
1199 M(tile_5d, "tile.5d") \
1200 M(tile_gather4_2d, "tile.gather4.2d") \
1201 M(im2col_3d, "im2col.3d") \
1202 M(im2col_4d, "im2col.4d") \
1203 M(im2col_5d, "im2col.5d") \
1204 M(im2col_w_3d, "im2col.w.3d") \
1205 M(im2col_w_4d, "im2col.w.4d") \
1206 M(im2col_w_5d, "im2col.w.5d") \
1207 M(im2col_w_128_3d, "im2col.w.128.3d") \
1208 M(im2col_w_128_4d, "im2col.w.128.4d") \
1209 M(im2col_w_128_5d, "im2col.w.128.5d")
1221 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s."))
1224#define G2S_ID(ID_SUFFIX, NAME) \
1225 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX)
1235 size_t NumParams =
F->getFunctionType()->getNumParams();
1239 if (!
F->getFunctionType()->getParamType(NumParams - 2)->isIntegerTy(1))
1246 Params[NumParams - 1]->isIntegerTy(1) ? NumParams - 4 : NumParams - 5;
1247 assert(Params[MaskIdx + 1]->isIntegerTy(64) &&
1248 "expected the i64 cache-hint after the multicast mask");
1249 Type *MaskTy = Params[MaskIdx];
1264 if (!Name.consume_front(
"cp.async.bulk.tensor.g2s.cta."))
1267#define G2S_CTA_ID(ID_SUFFIX, NAME) \
1268 .Case(NAME, Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX)
1280 if (!
F->getFunctionType()
1281 ->getParamType(
F->getFunctionType()->getNumParams() - 1)
1297 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cluster"))
1302 size_t NumParams =
F->getFunctionType()->getNumParams();
1303 if (!
F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
1307 Type *MaskTy =
F->getFunctionType()->getParamType(NumParams - 4);
1312 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
1325 if (!Name.consume_front(
"cp.async.bulk.global.to.shared.cta"))
1330 if (!
F->getFunctionType()->getParamType(5)->isIntegerTy(1))
1333 return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta;
1353 if (!Name.consume_front(
"cp.async.bulk.tensor.reduce."))
1356 auto [RedOpName, ShapeName] = Name.split(
'.');
1361 .
Case(
"tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
1362 .
Case(
"tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
1363 .
Case(
"tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
1364 .
Case(
"tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
1365 .
Case(
"tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
1366 .
Case(
"im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
1367 .
Case(
"im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
1368 .
Case(
"im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
1374 if (Name.consume_front(
"mapa.shared.cluster"))
1375 if (
F->getReturnType()->getPointerAddressSpace() ==
1377 return Intrinsic::nvvm_mapa_shared_cluster;
1379 if (Name.consume_front(
"cp.async.bulk.")) {
1382 .
Case(
"shared.cta.to.cluster",
1383 Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster)
1387 if (
F->getArg(0)->getType()->getPointerAddressSpace() ==
1397 if (!Name.consume_front(
"tcgen05.commit."))
1400 if (Name.consume_front(
"shared."))
1402 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_cg1)
1403 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_cg2)
1406 if (Name.consume_front(
"mc.shared.")) {
1408 if (!
F->getArg(1)->getType()->isIntegerTy(16))
1412 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_commit_mc_cg1)
1413 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_commit_mc_cg2)
1422 if (
F->arg_size() != 2)
1425 if (Name.consume_front(
"tcgen05.alloc.shared.") ||
1426 Name.consume_front(
"tcgen05.alloc."))
1428 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_alloc_cg1)
1429 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_alloc_cg2)
1432 if (Name.consume_front(
"tcgen05.dealloc."))
1434 .
Case(
"cg1", Intrinsic::nvvm_tcgen05_dealloc_cg1)
1435 .
Case(
"cg2", Intrinsic::nvvm_tcgen05_dealloc_cg2)
1442 if (Name.consume_front(
"fmax."))
1444 .
Case(
"bf16", Intrinsic::nvvm_fmax_bf16)
1445 .
Case(
"bf16x2", Intrinsic::nvvm_fmax_bf16x2)
1446 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmax_ftz_bf16)
1447 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmax_ftz_bf16x2)
1448 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmax_ftz_nan_bf16)
1449 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmax_ftz_nan_bf16x2)
1450 .
Case(
"ftz.nan.xorsign.abs.bf16",
1451 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16)
1452 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1453 Intrinsic::nvvm_fmax_ftz_nan_xorsign_abs_bf16x2)
1454 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16)
1455 .
Case(
"ftz.xorsign.abs.bf16x2",
1456 Intrinsic::nvvm_fmax_ftz_xorsign_abs_bf16x2)
1457 .
Case(
"nan.bf16", Intrinsic::nvvm_fmax_nan_bf16)
1458 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmax_nan_bf16x2)
1459 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16)
1460 .
Case(
"nan.xorsign.abs.bf16x2",
1461 Intrinsic::nvvm_fmax_nan_xorsign_abs_bf16x2)
1462 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmax_xorsign_abs_bf16)
1463 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmax_xorsign_abs_bf16x2)
1466 if (Name.consume_front(
"fmin."))
1468 .
Case(
"bf16", Intrinsic::nvvm_fmin_bf16)
1469 .
Case(
"bf16x2", Intrinsic::nvvm_fmin_bf16x2)
1470 .
Case(
"ftz.bf16", Intrinsic::nvvm_fmin_ftz_bf16)
1471 .
Case(
"ftz.bf16x2", Intrinsic::nvvm_fmin_ftz_bf16x2)
1472 .
Case(
"ftz.nan.bf16", Intrinsic::nvvm_fmin_ftz_nan_bf16)
1473 .
Case(
"ftz.nan.bf16x2", Intrinsic::nvvm_fmin_ftz_nan_bf16x2)
1474 .
Case(
"ftz.nan.xorsign.abs.bf16",
1475 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16)
1476 .
Case(
"ftz.nan.xorsign.abs.bf16x2",
1477 Intrinsic::nvvm_fmin_ftz_nan_xorsign_abs_bf16x2)
1478 .
Case(
"ftz.xorsign.abs.bf16", Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16)
1479 .
Case(
"ftz.xorsign.abs.bf16x2",
1480 Intrinsic::nvvm_fmin_ftz_xorsign_abs_bf16x2)
1481 .
Case(
"nan.bf16", Intrinsic::nvvm_fmin_nan_bf16)
1482 .
Case(
"nan.bf16x2", Intrinsic::nvvm_fmin_nan_bf16x2)
1483 .
Case(
"nan.xorsign.abs.bf16", Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16)
1484 .
Case(
"nan.xorsign.abs.bf16x2",
1485 Intrinsic::nvvm_fmin_nan_xorsign_abs_bf16x2)
1486 .
Case(
"xorsign.abs.bf16", Intrinsic::nvvm_fmin_xorsign_abs_bf16)
1487 .
Case(
"xorsign.abs.bf16x2", Intrinsic::nvvm_fmin_xorsign_abs_bf16x2)
1490 if (Name.consume_front(
"neg."))
1492 .
Case(
"bf16", Intrinsic::nvvm_neg_bf16)
1493 .
Case(
"bf16x2", Intrinsic::nvvm_neg_bf16x2)
1502 auto IsOldBF16StorageTy = [](
Type *OldTy,
Type *NewTy) {
1507 if (!IsOldBF16StorageTy(OldFnTy->getReturnType(), NewFnTy->getReturnType()))
1510 if (OldFnTy->getNumParams() != NewFnTy->getNumParams())
1513 for (
unsigned I = 0,
E = OldFnTy->getNumParams();
I !=
E; ++
I)
1514 if (!IsOldBF16StorageTy(OldFnTy->getParamType(
I), NewFnTy->getParamType(
I)))
1524 {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat,
1528 {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat,
1531 {Intrinsic::nvvm_fma, Intrinsic::nvvm_fma_sat, Intrinsic::nvvm_fma_relu},
1532 {Intrinsic::nvvm_fma_ftz, Intrinsic::nvvm_fma_ftz_sat,
1533 Intrinsic::nvvm_fma_ftz_relu}};
1539 bool IsFTZ,
bool IsSat,
1544 .
Case(
"f", !IsRelu && !(IsSat && IsFMul))
1545 .
Case(
"d", !IsFTZ && !IsSat && !IsRelu)
1546 .
Case(
"f16", IsRN && (IsFMA || IsSat))
1547 .
Case(
"v2f16", IsRN && !IsFMA && IsSat)
1548 .
Case(
"f16x2", IsRN && IsFMA)
1549 .
Cases({
"bf16",
"bf16x2"}, IsRN && IsFMA && !IsFTZ && !IsSat)
1553static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
1555 auto [Modifiers,
Type] = Name.rsplit(
'.');
1565 return std::nullopt;
1567 StringRef Rest = Modifiers.drop_front(2);
1572 return std::nullopt;
1574 const Intrinsic::ID IID = IIDs[IsFTZ][IsSat ? 1 : IsRelu ? 2 : 0];
1576 return std::nullopt;
1582 return std::nullopt;
1588 if (Name !=
"mbarrier.init" && Name !=
"mbarrier.init.shared")
1591 return Intrinsic::nvvm_mbarrier_init;
1595 return Name.consume_front(
"local") || Name.consume_front(
"shared") ||
1596 Name.consume_front(
"global") || Name.consume_front(
"constant") ||
1597 Name.consume_front(
"param");
1601 if (!Name.consume_front(
"vp."))
1630 .
StartsWith(
"ptrtoint", Instruction::PtrToInt)
1631 .
StartsWith(
"inttoptr", Instruction::IntToPtr)
1638 if (!Name.consume_front(
"vp."))
1658 .
StartsWith(
"nearbyint", Intrinsic::nearbyint)
1659 .
StartsWith(
"roundeven", Intrinsic::roundeven)
1664 .
StartsWith(
"bitreverse", Intrinsic::bitreverse)
1676 .
StartsWith(
"is.fpclass", Intrinsic::is_fpclass)
1687 if (Name.starts_with(
"to.fp16")) {
1691 FuncTy->getReturnType());
1694 if (Name.starts_with(
"from.fp16")) {
1698 FuncTy->getReturnType());
1708 if (Defaults.empty())
1711 unsigned FullArgCount = FirstDefault + Defaults.size();
1714 if (
F->arg_size() < FirstDefault ||
F->arg_size() >= FullArgCount)
1717 unsigned NumMissingTrailingParams = FullArgCount -
F->arg_size();
1719 NumMissingTrailingParams))
1722 return FullArgCount;
1729 unsigned FullArgCount =
1731 if (FullArgCount == 0)
1737 "total number of default args does not match intrinsic signature");
1742 bool CanUpgradeDebugIntrinsicsToRecords) {
1743 assert(
F &&
"Illegal to upgrade a non-existent Function.");
1748 if (!Name.consume_front(
"llvm.") || Name.empty())
1754 bool IsArm = Name.consume_front(
"arm.");
1755 if (IsArm || Name.consume_front(
"aarch64.")) {
1761 if (Name.consume_front(
"amdgcn.")) {
1762 if (Name ==
"alignbit") {
1765 F->getParent(), Intrinsic::fshr, {F->getReturnType()});
1769 if (Name.consume_front(
"atomic.")) {
1770 if (Name.starts_with(
"inc") || Name.starts_with(
"dec") ||
1771 Name.starts_with(
"cond.sub") || Name.starts_with(
"csub")) {
1780 if (Name.starts_with(
"addrspacecast.nonnull")) {
1787 switch (
F->getIntrinsicID()) {
1791 case Intrinsic::amdgcn_wmma_i32_16x16x64_iu8:
1792 if (
F->arg_size() == 7) {
1797 case Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8:
1798 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
1799 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
1800 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
1801 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
1802 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
1803 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16:
1804 if (
F->arg_size() == 8) {
1811 if (Name.consume_front(
"ds.") || Name.consume_front(
"global.atomic.") ||
1812 Name.consume_front(
"flat.atomic.")) {
1813 if (Name.starts_with(
"fadd") ||
1815 (Name.starts_with(
"fmin") && !Name.starts_with(
"fmin.num")) ||
1816 (Name.starts_with(
"fmax") && !Name.starts_with(
"fmax.num"))) {
1824 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
1829 if (Name.starts_with(
"ldexp.")) {
1832 F->getParent(), Intrinsic::ldexp,
1833 {F->getReturnType(), F->getArg(1)->getType()});
1842 if (
F->arg_size() == 1) {
1843 if (Name.consume_front(
"convert.")) {
1857 F->arg_begin()->getType());
1863 if (Name ==
"coro.end" &&
1864 (
F->arg_size() == 2 ||
F->getReturnType()->isIntegerTy(1)))
1865 CoroEndID = Intrinsic::coro_end;
1866 else if (Name ==
"coro.end.async" &&
F->getReturnType()->isIntegerTy(1))
1867 CoroEndID = Intrinsic::coro_end_async;
1878 if (Name.consume_front(
"dbg.")) {
1880 if (CanUpgradeDebugIntrinsicsToRecords) {
1881 if (Name ==
"addr" || Name ==
"value" || Name ==
"assign" ||
1882 Name ==
"declare" || Name ==
"label") {
1891 if (Name ==
"addr" || (Name ==
"value" &&
F->arg_size() == 4)) {
1894 Intrinsic::dbg_value);
1901 if (Name.consume_front(
"experimental.vector.")) {
1907 .
StartsWith(
"extract.", Intrinsic::vector_extract)
1908 .
StartsWith(
"insert.", Intrinsic::vector_insert)
1909 .
StartsWith(
"reverse.", Intrinsic::vector_reverse)
1910 .
StartsWith(
"interleave2.", Intrinsic::vector_interleave2)
1911 .
StartsWith(
"deinterleave2.", Intrinsic::vector_deinterleave2)
1913 Intrinsic::vector_partial_reduce_add)
1916 const auto *FT =
F->getFunctionType();
1918 if (ID == Intrinsic::vector_extract ||
1919 ID == Intrinsic::vector_interleave2)
1922 if (ID != Intrinsic::vector_interleave2)
1924 if (ID == Intrinsic::vector_insert ||
1925 ID == Intrinsic::vector_partial_reduce_add)
1933 if (Name.consume_front(
"reduce.")) {
1935 static const Regex R(
"^([a-z]+)\\.[a-z][0-9]+");
1936 if (R.match(Name, &
Groups))
1938 .
Case(
"add", Intrinsic::vector_reduce_add)
1939 .
Case(
"mul", Intrinsic::vector_reduce_mul)
1940 .
Case(
"and", Intrinsic::vector_reduce_and)
1941 .
Case(
"or", Intrinsic::vector_reduce_or)
1942 .
Case(
"xor", Intrinsic::vector_reduce_xor)
1943 .
Case(
"smax", Intrinsic::vector_reduce_smax)
1944 .
Case(
"smin", Intrinsic::vector_reduce_smin)
1945 .
Case(
"umax", Intrinsic::vector_reduce_umax)
1946 .
Case(
"umin", Intrinsic::vector_reduce_umin)
1947 .
Case(
"fmax", Intrinsic::vector_reduce_fmax)
1948 .
Case(
"fmin", Intrinsic::vector_reduce_fmin)
1953 static const Regex R2(
"^v2\\.([a-z]+)\\.[fi][0-9]+");
1958 .
Case(
"fadd", Intrinsic::vector_reduce_fadd)
1959 .
Case(
"fmul", Intrinsic::vector_reduce_fmul)
1964 auto Args =
F->getFunctionType()->params();
1966 {Args[V2 ? 1 : 0]});
1972 if (Name.consume_front(
"splice"))
1976 if (Name.consume_front(
"experimental.stepvector.")) {
1980 F->getParent(), ID,
F->getFunctionType()->getReturnType());
1985 if (Name.starts_with(
"flt.rounds")) {
1988 Intrinsic::get_rounding);
1993 if (Name.starts_with(
"invariant.group.barrier")) {
1995 auto Args =
F->getFunctionType()->params();
1996 Type* ObjectPtr[1] = {Args[0]};
1999 F->getParent(), Intrinsic::launder_invariant_group, ObjectPtr);
2004 bool IsLifetimeStart = Name.consume_front(
"lifetime.start");
2005 bool IsLifetimeEnd = !IsLifetimeStart && Name.consume_front(
"lifetime.end");
2006 if (IsLifetimeStart || IsLifetimeEnd) {
2007 if (
F->arg_size() == 2) {
2008 Intrinsic::ID IID = IsLifetimeStart ? Intrinsic::lifetime_start
2009 : Intrinsic::lifetime_end;
2014 F->getArg(1)->getType());
2016 }
else if (
F->arg_size() == 1 && Name ==
".i64") {
2036 .StartsWith(
"memcpy.", Intrinsic::memcpy)
2037 .StartsWith(
"memmove.", Intrinsic::memmove)
2039 if (
F->arg_size() == 5) {
2043 F->getFunctionType()->params().slice(0, 3);
2049 if (Name.starts_with(
"memset.") &&
F->arg_size() == 5) {
2052 const auto *FT =
F->getFunctionType();
2053 Type *ParamTypes[2] = {
2054 FT->getParamType(0),
2058 Intrinsic::memset, ParamTypes);
2064 .
StartsWith(
"masked.load", Intrinsic::masked_load)
2065 .
StartsWith(
"masked.gather", Intrinsic::masked_gather)
2066 .
StartsWith(
"masked.store", Intrinsic::masked_store)
2067 .
StartsWith(
"masked.scatter", Intrinsic::masked_scatter)
2069 if (MaskedID &&
F->arg_size() == 4) {
2071 if (MaskedID == Intrinsic::masked_load ||
2072 MaskedID == Intrinsic::masked_gather) {
2074 F->getParent(), MaskedID,
2075 {F->getReturnType(), F->getArg(0)->getType()});
2079 F->getParent(), MaskedID,
2080 {F->getArg(0)->getType(), F->getArg(1)->getType()});
2086 if (Name.consume_front(
"nvvm.")) {
2088 if (
F->arg_size() == 1) {
2091 .
Cases({
"brev32",
"brev64"}, Intrinsic::bitreverse)
2092 .Case(
"clz.i", Intrinsic::ctlz)
2093 .
Case(
"popc.i", Intrinsic::ctpop)
2097 {F->getReturnType()});
2100 }
else if (
F->arg_size() == 2) {
2103 .
Cases({
"max.s",
"max.i",
"max.ll"}, Intrinsic::smax)
2104 .Cases({
"min.s",
"min.i",
"min.ll"}, Intrinsic::smin)
2105 .Cases({
"max.us",
"max.ui",
"max.ull"}, Intrinsic::umax)
2106 .Cases({
"min.us",
"min.ui",
"min.ull"}, Intrinsic::umin)
2107 .Cases({
"mulhi.s",
"mulhi.i",
"mulhi.ll"}, Intrinsic::smulh)
2108 .Cases({
"mulhi.us",
"mulhi.ui",
"mulhi.ull"}, Intrinsic::umulh)
2112 {F->getReturnType()});
2149 F->getParent(), IID,
F->getReturnType(),
2150 F->getFunctionType()->params());
2161 {F->getArg(0)->getType()});
2209 F->getArg(0)->getType());
2217 bool Expand =
false;
2218 if (Name.consume_front(
"abs."))
2221 Name ==
"i" || Name ==
"ll" || Name ==
"bf16" || Name ==
"bf16x2";
2222 else if (Name.consume_front(
"fabs."))
2224 Expand = Name ==
"f" || Name ==
"ftz.f" || Name ==
"d";
2225 else if (Name.consume_front(
"add."))
2228 else if (Name.consume_front(
"mul."))
2231 else if (Name.consume_front(
"ex2.approx."))
2234 Name ==
"f" || Name ==
"ftz.f" || Name ==
"d" || Name ==
"f16x2";
2235 else if (Name.consume_front(
"fma."))
2238 else if (Name.consume_front(
"atomic.load."))
2247 else if (Name.consume_front(
"atomic."))
2262 else if (Name.consume_front(
"bitcast."))
2265 Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" || Name ==
"d2ll";
2266 else if (Name.consume_front(
"rotate."))
2268 Expand = Name ==
"b32" || Name ==
"b64" || Name ==
"right.b64";
2269 else if (Name.consume_front(
"ptr.gen.to."))
2272 else if (Name.consume_front(
"ptr."))
2275 else if (Name.consume_front(
"ldg.global."))
2277 Expand = (Name.starts_with(
"i.") || Name.starts_with(
"f.") ||
2278 Name.starts_with(
"p."));
2281 .
Case(
"barrier0",
true)
2282 .
Case(
"barrier.n",
true)
2283 .
Case(
"barrier.sync.cnt",
true)
2284 .
Case(
"barrier.sync",
true)
2285 .
Case(
"barrier",
true)
2286 .
Case(
"bar.sync",
true)
2287 .
Case(
"barrier0.popc",
true)
2288 .
Case(
"barrier0.and",
true)
2289 .
Case(
"barrier0.or",
true)
2290 .
Case(
"clz.ll",
true)
2291 .
Case(
"popc.ll",
true)
2293 .
Case(
"swap.lo.hi.b64",
true)
2294 .
Case(
"tanh.approx.f32",
true)
2306 if (Name.starts_with(
"objectsize.")) {
2307 Type *Tys[2] = {
F->getReturnType(),
F->arg_begin()->getType() };
2308 if (
F->arg_size() == 2 ||
F->arg_size() == 3) {
2311 Intrinsic::objectsize, Tys);
2318 if (Name.starts_with(
"ptr.annotation.") &&
F->arg_size() == 4) {
2321 F->getParent(), Intrinsic::ptr_annotation,
2322 {F->arg_begin()->getType(), F->getArg(1)->getType()});
2328 if (Name.consume_front(
"riscv.")) {
2331 .
Case(
"aes32dsi", Intrinsic::riscv_aes32dsi)
2332 .
Case(
"aes32dsmi", Intrinsic::riscv_aes32dsmi)
2333 .
Case(
"aes32esi", Intrinsic::riscv_aes32esi)
2334 .
Case(
"aes32esmi", Intrinsic::riscv_aes32esmi)
2337 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32)) {
2350 if (!
F->getFunctionType()->getParamType(2)->isIntegerTy(32) ||
2351 F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2360 .
StartsWith(
"sha256sig0", Intrinsic::riscv_sha256sig0)
2361 .
StartsWith(
"sha256sig1", Intrinsic::riscv_sha256sig1)
2362 .
StartsWith(
"sha256sum0", Intrinsic::riscv_sha256sum0)
2363 .
StartsWith(
"sha256sum1", Intrinsic::riscv_sha256sum1)
2368 if (
F->getFunctionType()->getReturnType()->isIntegerTy(64)) {
2377 if (Name ==
"clmul.i32" || Name ==
"clmul.i64") {
2379 F->getParent(), Intrinsic::clmul, {F->getReturnType()});
2388 if (Name ==
"stackprotectorcheck") {
2392 if (Name.starts_with(
"strip.invariant.group")) {
2397 F->getParent(), Intrinsic::launder_invariant_group,
2398 F->getReturnType());
2404 if (Name ==
"thread.pointer") {
2406 F->getParent(), Intrinsic::thread_pointer,
F->getReturnType());
2412 if (Name ==
"var.annotation" &&
F->arg_size() == 4) {
2415 F->getParent(), Intrinsic::var_annotation,
2416 {{F->arg_begin()->getType(), F->getArg(1)->getType()}});
2419 if (Name.consume_front(
"vector.splice")) {
2420 if (Name.starts_with(
".left") || Name.starts_with(
".right"))
2430 if (Name.consume_front(
"wasm.")) {
2433 .
StartsWith(
"fma.", Intrinsic::wasm_relaxed_madd)
2434 .
StartsWith(
"fms.", Intrinsic::wasm_relaxed_nmadd)
2435 .
StartsWith(
"laneselect.", Intrinsic::wasm_relaxed_laneselect)
2440 F->getReturnType());
2444 if (Name.consume_front(
"dot.i8x16.i7x16.")) {
2446 .
Case(
"signed", Intrinsic::wasm_relaxed_dot_i8x16_i7x16_signed)
2448 Intrinsic::wasm_relaxed_dot_i8x16_i7x16_add_signed)
2467 if (ST && (!
ST->isLiteral() ||
ST->isPacked()) &&
2477 std::string
Name =
F->getName().str();
2480 Name,
F->getParent());
2491 if (Result != std::nullopt) {
2508 bool CanUpgradeDebugIntrinsicsToRecords) {
2528 GV->
getName() ==
"llvm.global_dtors")) ||
2542 unsigned N =
Init->getNumOperands();
2543 std::vector<Constant *> NewCtors(
N);
2544 for (
unsigned i = 0; i !=
N; ++i) {
2547 Ctor->getAggregateElement(1),
2561 unsigned NumElts = ResultTy->getNumElements() * 8;
2565 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2575 for (
unsigned l = 0; l != NumElts; l += 16)
2576 for (
unsigned i = 0; i != 16; ++i) {
2577 unsigned Idx = NumElts + i - Shift;
2579 Idx -= NumElts - 16;
2580 Idxs[l + i] = Idx + l;
2583 Res = Builder.CreateShuffleVector(Res,
Op,
ArrayRef(Idxs, NumElts));
2587 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2595 unsigned NumElts = ResultTy->getNumElements() * 8;
2599 Op = Builder.CreateBitCast(
Op, VecTy,
"cast");
2609 for (
unsigned l = 0; l != NumElts; l += 16)
2610 for (
unsigned i = 0; i != 16; ++i) {
2611 unsigned Idx = i + Shift;
2613 Idx += NumElts - 16;
2614 Idxs[l + i] = Idx + l;
2617 Res = Builder.CreateShuffleVector(
Op, Res,
ArrayRef(Idxs, NumElts));
2621 return Builder.CreateBitCast(Res, ResultTy,
"cast");
2629 Mask = Builder.CreateBitCast(Mask, MaskTy);
2635 for (
unsigned i = 0; i != NumElts; ++i)
2637 Mask = Builder.CreateShuffleVector(Mask, Mask,
ArrayRef(Indices, NumElts),
2648 if (
C->isAllOnesValue())
2653 return Builder.CreateSelect(Mask, Op0, Op1);
2660 if (
C->isAllOnesValue())
2664 Mask->getType()->getIntegerBitWidth());
2665 Mask = Builder.CreateBitCast(Mask, MaskTy);
2666 Mask = Builder.CreateExtractElement(Mask, (
uint64_t)0);
2667 return Builder.CreateSelect(Mask, Op0, Op1);
2680 assert((IsVALIGN || NumElts % 16 == 0) &&
"Illegal NumElts for PALIGNR!");
2681 assert((!IsVALIGN || NumElts <= 16) &&
"NumElts too large for VALIGN!");
2686 ShiftVal &= (NumElts - 1);
2695 if (ShiftVal > 16) {
2703 for (
unsigned l = 0; l < NumElts; l += 16) {
2704 for (
unsigned i = 0; i != 16; ++i) {
2705 unsigned Idx = ShiftVal + i;
2706 if (!IsVALIGN && Idx >= 16)
2707 Idx += NumElts - 16;
2708 Indices[l + i] = Idx + l;
2713 Op1, Op0,
ArrayRef(Indices, NumElts),
"palignr");
2719 bool ZeroMask,
bool IndexForm) {
2722 unsigned EltWidth = Ty->getScalarSizeInBits();
2723 bool IsFloat = Ty->isFPOrFPVectorTy();
2725 if (VecWidth == 128 && EltWidth == 32 && IsFloat)
2726 IID = Intrinsic::x86_avx512_vpermi2var_ps_128;
2727 else if (VecWidth == 128 && EltWidth == 32 && !IsFloat)
2728 IID = Intrinsic::x86_avx512_vpermi2var_d_128;
2729 else if (VecWidth == 128 && EltWidth == 64 && IsFloat)
2730 IID = Intrinsic::x86_avx512_vpermi2var_pd_128;
2731 else if (VecWidth == 128 && EltWidth == 64 && !IsFloat)
2732 IID = Intrinsic::x86_avx512_vpermi2var_q_128;
2733 else if (VecWidth == 256 && EltWidth == 32 && IsFloat)
2734 IID = Intrinsic::x86_avx512_vpermi2var_ps_256;
2735 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
2736 IID = Intrinsic::x86_avx512_vpermi2var_d_256;
2737 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
2738 IID = Intrinsic::x86_avx512_vpermi2var_pd_256;
2739 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
2740 IID = Intrinsic::x86_avx512_vpermi2var_q_256;
2741 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
2742 IID = Intrinsic::x86_avx512_vpermi2var_ps_512;
2743 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
2744 IID = Intrinsic::x86_avx512_vpermi2var_d_512;
2745 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
2746 IID = Intrinsic::x86_avx512_vpermi2var_pd_512;
2747 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
2748 IID = Intrinsic::x86_avx512_vpermi2var_q_512;
2749 else if (VecWidth == 128 && EltWidth == 16)
2750 IID = Intrinsic::x86_avx512_vpermi2var_hi_128;
2751 else if (VecWidth == 256 && EltWidth == 16)
2752 IID = Intrinsic::x86_avx512_vpermi2var_hi_256;
2753 else if (VecWidth == 512 && EltWidth == 16)
2754 IID = Intrinsic::x86_avx512_vpermi2var_hi_512;
2755 else if (VecWidth == 128 && EltWidth == 8)
2756 IID = Intrinsic::x86_avx512_vpermi2var_qi_128;
2757 else if (VecWidth == 256 && EltWidth == 8)
2758 IID = Intrinsic::x86_avx512_vpermi2var_qi_256;
2759 else if (VecWidth == 512 && EltWidth == 8)
2760 IID = Intrinsic::x86_avx512_vpermi2var_qi_512;
2771 Value *V = Builder.CreateIntrinsic(IID, Args);
2783 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1});
2794 bool IsRotateRight) {
2804 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2805 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2808 Intrinsic::ID IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
2809 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Src, Src, Amt});
2854 Value *Ext = Builder.CreateSExt(Cmp, Ty);
2859 bool IsShiftRight,
bool ZeroMask) {
2873 Amt = Builder.CreateIntCast(Amt, Ty->getScalarType(),
false);
2874 Amt = Builder.CreateVectorSplat(NumElts, Amt);
2877 Intrinsic::ID IID = IsShiftRight ? Intrinsic::fshr : Intrinsic::fshl;
2878 Value *Res = Builder.CreateIntrinsic(IID, Ty, {Op0, Op1, Amt});
2893 const Align Alignment =
2895 ?
Align(
Data->getType()->getPrimitiveSizeInBits().getFixedValue() / 8)
2900 if (
C->isAllOnesValue())
2901 return Builder.CreateAlignedStore(
Data, Ptr, Alignment);
2906 return Builder.CreateMaskedStore(
Data, Ptr, Alignment, Mask);
2912 const Align Alignment =
2921 if (
C->isAllOnesValue())
2922 return Builder.CreateAlignedLoad(ValTy, Ptr, Alignment);
2927 return Builder.CreateMaskedLoad(ValTy, Ptr, Alignment, Mask, Passthru);
2933 Value *Res = Builder.CreateIntrinsic(Intrinsic::abs, Ty,
2934 {Op0, Builder.getInt1(
false)});
2949 Constant *ShiftAmt = ConstantInt::get(Ty, 32);
2950 LHS = Builder.CreateShl(
LHS, ShiftAmt);
2951 LHS = Builder.CreateAShr(
LHS, ShiftAmt);
2952 RHS = Builder.CreateShl(
RHS, ShiftAmt);
2953 RHS = Builder.CreateAShr(
RHS, ShiftAmt);
2956 Constant *Mask = ConstantInt::get(Ty, 0xffffffff);
2957 LHS = Builder.CreateAnd(
LHS, Mask);
2958 RHS = Builder.CreateAnd(
RHS, Mask);
2975 if (!
C || !
C->isAllOnesValue())
2976 Vec = Builder.CreateAnd(Vec,
getX86MaskVec(Builder, Mask, NumElts));
2981 for (
unsigned i = 0; i != NumElts; ++i)
2983 for (
unsigned i = NumElts; i != 8; ++i)
2984 Indices[i] = NumElts + i % NumElts;
2985 Vec = Builder.CreateShuffleVector(Vec,
2989 return Builder.CreateBitCast(Vec, Builder.getIntNTy(std::max(NumElts, 8U)));
3001 }
else if (
CC == 7) {
3037 Value* AndNode = Builder.CreateAnd(Mask,
APInt(8, 1));
3038 Value* Cmp = Builder.CreateIsNotNull(AndNode);
3040 Value* Extract2 = Builder.CreateExtractElement(Src, (
uint64_t)0);
3041 Value*
Select = Builder.CreateSelect(Cmp, Extract1, Extract2);
3050 return Builder.CreateSExt(Mask, ReturnOp,
"vpmovm2");
3056 Name = Name.substr(12);
3061 if (Name.starts_with(
"max.p")) {
3062 if (VecWidth == 128 && EltWidth == 32)
3063 IID = Intrinsic::x86_sse_max_ps;
3064 else if (VecWidth == 128 && EltWidth == 64)
3065 IID = Intrinsic::x86_sse2_max_pd;
3066 else if (VecWidth == 256 && EltWidth == 32)
3067 IID = Intrinsic::x86_avx_max_ps_256;
3068 else if (VecWidth == 256 && EltWidth == 64)
3069 IID = Intrinsic::x86_avx_max_pd_256;
3072 }
else if (Name.starts_with(
"min.p")) {
3073 if (VecWidth == 128 && EltWidth == 32)
3074 IID = Intrinsic::x86_sse_min_ps;
3075 else if (VecWidth == 128 && EltWidth == 64)
3076 IID = Intrinsic::x86_sse2_min_pd;
3077 else if (VecWidth == 256 && EltWidth == 32)
3078 IID = Intrinsic::x86_avx_min_ps_256;
3079 else if (VecWidth == 256 && EltWidth == 64)
3080 IID = Intrinsic::x86_avx_min_pd_256;
3083 }
else if (Name.starts_with(
"pshuf.b.")) {
3084 if (VecWidth == 128)
3085 IID = Intrinsic::x86_ssse3_pshuf_b_128;
3086 else if (VecWidth == 256)
3087 IID = Intrinsic::x86_avx2_pshuf_b;
3088 else if (VecWidth == 512)
3089 IID = Intrinsic::x86_avx512_pshuf_b_512;
3092 }
else if (Name.starts_with(
"pmul.hr.sw.")) {
3093 if (VecWidth == 128)
3094 IID = Intrinsic::x86_ssse3_pmul_hr_sw_128;
3095 else if (VecWidth == 256)
3096 IID = Intrinsic::x86_avx2_pmul_hr_sw;
3097 else if (VecWidth == 512)
3098 IID = Intrinsic::x86_avx512_pmul_hr_sw_512;
3101 }
else if (Name.starts_with(
"pmulh.w")) {
3102 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3103 "Unexpected intrinsic");
3106 }
else if (Name.starts_with(
"pmulhu.w")) {
3107 assert((VecWidth == 128 || VecWidth == 256 || VecWidth == 512) &&
3108 "Unexpected intrinsic");
3111 }
else if (Name.starts_with(
"pmaddw.d.")) {
3112 if (VecWidth == 128)
3113 IID = Intrinsic::x86_sse2_pmadd_wd;
3114 else if (VecWidth == 256)
3115 IID = Intrinsic::x86_avx2_pmadd_wd;
3116 else if (VecWidth == 512)
3117 IID = Intrinsic::x86_avx512_pmaddw_d_512;
3120 }
else if (Name.starts_with(
"pmaddubs.w.")) {
3121 if (VecWidth == 128)
3122 IID = Intrinsic::x86_ssse3_pmadd_ub_sw_128;
3123 else if (VecWidth == 256)
3124 IID = Intrinsic::x86_avx2_pmadd_ub_sw;
3125 else if (VecWidth == 512)
3126 IID = Intrinsic::x86_avx512_pmaddubs_w_512;
3129 }
else if (Name.starts_with(
"packsswb.")) {
3130 if (VecWidth == 128)
3131 IID = Intrinsic::x86_sse2_packsswb_128;
3132 else if (VecWidth == 256)
3133 IID = Intrinsic::x86_avx2_packsswb;
3134 else if (VecWidth == 512)
3135 IID = Intrinsic::x86_avx512_packsswb_512;
3138 }
else if (Name.starts_with(
"packssdw.")) {
3139 if (VecWidth == 128)
3140 IID = Intrinsic::x86_sse2_packssdw_128;
3141 else if (VecWidth == 256)
3142 IID = Intrinsic::x86_avx2_packssdw;
3143 else if (VecWidth == 512)
3144 IID = Intrinsic::x86_avx512_packssdw_512;
3147 }
else if (Name.starts_with(
"packuswb.")) {
3148 if (VecWidth == 128)
3149 IID = Intrinsic::x86_sse2_packuswb_128;
3150 else if (VecWidth == 256)
3151 IID = Intrinsic::x86_avx2_packuswb;
3152 else if (VecWidth == 512)
3153 IID = Intrinsic::x86_avx512_packuswb_512;
3156 }
else if (Name.starts_with(
"packusdw.")) {
3157 if (VecWidth == 128)
3158 IID = Intrinsic::x86_sse41_packusdw;
3159 else if (VecWidth == 256)
3160 IID = Intrinsic::x86_avx2_packusdw;
3161 else if (VecWidth == 512)
3162 IID = Intrinsic::x86_avx512_packusdw_512;
3165 }
else if (Name.starts_with(
"vpermilvar.")) {
3166 if (VecWidth == 128 && EltWidth == 32)
3167 IID = Intrinsic::x86_avx_vpermilvar_ps;
3168 else if (VecWidth == 128 && EltWidth == 64)
3169 IID = Intrinsic::x86_avx_vpermilvar_pd;
3170 else if (VecWidth == 256 && EltWidth == 32)
3171 IID = Intrinsic::x86_avx_vpermilvar_ps_256;
3172 else if (VecWidth == 256 && EltWidth == 64)
3173 IID = Intrinsic::x86_avx_vpermilvar_pd_256;
3174 else if (VecWidth == 512 && EltWidth == 32)
3175 IID = Intrinsic::x86_avx512_vpermilvar_ps_512;
3176 else if (VecWidth == 512 && EltWidth == 64)
3177 IID = Intrinsic::x86_avx512_vpermilvar_pd_512;
3180 }
else if (Name ==
"cvtpd2dq.256") {
3181 IID = Intrinsic::x86_avx_cvt_pd2dq_256;
3182 }
else if (Name ==
"cvtpd2ps.256") {
3183 IID = Intrinsic::x86_avx_cvt_pd2_ps_256;
3184 }
else if (Name ==
"cvttpd2dq.256") {
3185 IID = Intrinsic::x86_avx_cvtt_pd2dq_256;
3186 }
else if (Name ==
"cvttps2dq.128") {
3187 IID = Intrinsic::x86_sse2_cvttps2dq;
3188 }
else if (Name ==
"cvttps2dq.256") {
3189 IID = Intrinsic::x86_avx_cvtt_ps2dq_256;
3190 }
else if (Name.starts_with(
"permvar.")) {
3192 if (VecWidth == 256 && EltWidth == 32 && IsFloat)
3193 IID = Intrinsic::x86_avx2_permps;
3194 else if (VecWidth == 256 && EltWidth == 32 && !IsFloat)
3195 IID = Intrinsic::x86_avx2_permd;
3196 else if (VecWidth == 256 && EltWidth == 64 && IsFloat)
3197 IID = Intrinsic::x86_avx512_permvar_df_256;
3198 else if (VecWidth == 256 && EltWidth == 64 && !IsFloat)
3199 IID = Intrinsic::x86_avx512_permvar_di_256;
3200 else if (VecWidth == 512 && EltWidth == 32 && IsFloat)
3201 IID = Intrinsic::x86_avx512_permvar_sf_512;
3202 else if (VecWidth == 512 && EltWidth == 32 && !IsFloat)
3203 IID = Intrinsic::x86_avx512_permvar_si_512;
3204 else if (VecWidth == 512 && EltWidth == 64 && IsFloat)
3205 IID = Intrinsic::x86_avx512_permvar_df_512;
3206 else if (VecWidth == 512 && EltWidth == 64 && !IsFloat)
3207 IID = Intrinsic::x86_avx512_permvar_di_512;
3208 else if (VecWidth == 128 && EltWidth == 16)
3209 IID = Intrinsic::x86_avx512_permvar_hi_128;
3210 else if (VecWidth == 256 && EltWidth == 16)
3211 IID = Intrinsic::x86_avx512_permvar_hi_256;
3212 else if (VecWidth == 512 && EltWidth == 16)
3213 IID = Intrinsic::x86_avx512_permvar_hi_512;
3214 else if (VecWidth == 128 && EltWidth == 8)
3215 IID = Intrinsic::x86_avx512_permvar_qi_128;
3216 else if (VecWidth == 256 && EltWidth == 8)
3217 IID = Intrinsic::x86_avx512_permvar_qi_256;
3218 else if (VecWidth == 512 && EltWidth == 8)
3219 IID = Intrinsic::x86_avx512_permvar_qi_512;
3222 }
else if (Name.starts_with(
"dbpsadbw.")) {
3223 if (VecWidth == 128)
3224 IID = Intrinsic::x86_avx512_dbpsadbw_128;
3225 else if (VecWidth == 256)
3226 IID = Intrinsic::x86_avx512_dbpsadbw_256;
3227 else if (VecWidth == 512)
3228 IID = Intrinsic::x86_avx512_dbpsadbw_512;
3231 }
else if (Name.starts_with(
"pmultishift.qb.")) {
3232 if (VecWidth == 128)
3233 IID = Intrinsic::x86_avx512_pmultishift_qb_128;
3234 else if (VecWidth == 256)
3235 IID = Intrinsic::x86_avx512_pmultishift_qb_256;
3236 else if (VecWidth == 512)
3237 IID = Intrinsic::x86_avx512_pmultishift_qb_512;
3240 }
else if (Name.starts_with(
"conflict.")) {
3241 if (Name[9] ==
'd' && VecWidth == 128)
3242 IID = Intrinsic::x86_avx512_conflict_d_128;
3243 else if (Name[9] ==
'd' && VecWidth == 256)
3244 IID = Intrinsic::x86_avx512_conflict_d_256;
3245 else if (Name[9] ==
'd' && VecWidth == 512)
3246 IID = Intrinsic::x86_avx512_conflict_d_512;
3247 else if (Name[9] ==
'q' && VecWidth == 128)
3248 IID = Intrinsic::x86_avx512_conflict_q_128;
3249 else if (Name[9] ==
'q' && VecWidth == 256)
3250 IID = Intrinsic::x86_avx512_conflict_q_256;
3251 else if (Name[9] ==
'q' && VecWidth == 512)
3252 IID = Intrinsic::x86_avx512_conflict_q_512;
3255 }
else if (Name.starts_with(
"pavg.")) {
3256 if (Name[5] ==
'b' && VecWidth == 128)
3257 IID = Intrinsic::x86_sse2_pavg_b;
3258 else if (Name[5] ==
'b' && VecWidth == 256)
3259 IID = Intrinsic::x86_avx2_pavg_b;
3260 else if (Name[5] ==
'b' && VecWidth == 512)
3261 IID = Intrinsic::x86_avx512_pavg_b_512;
3262 else if (Name[5] ==
'w' && VecWidth == 128)
3263 IID = Intrinsic::x86_sse2_pavg_w;
3264 else if (Name[5] ==
'w' && VecWidth == 256)
3265 IID = Intrinsic::x86_avx2_pavg_w;
3266 else if (Name[5] ==
'w' && VecWidth == 512)
3267 IID = Intrinsic::x86_avx512_pavg_w_512;
3276 Rep = Builder.CreateIntrinsic(IID, Args);
3287 if (AsmStr->find(
"mov\tfp") == 0 &&
3288 AsmStr->find(
"objc_retainAutoreleaseReturnValue") != std::string::npos &&
3289 (Pos = AsmStr->find(
"# marker")) != std::string::npos) {
3290 AsmStr->replace(Pos, 1,
";");
3298 assert(Result &&
"unsupported nvvm.{add,mul,fma}.* intrinsic");
3306 if (IsLegacyBF16Storage)
3307 Ty = Name.ends_with(
"bf16x2")
3309 :
static_cast<Type *
>(Builder.getBFloatTy());
3313 Args.push_back(IsLegacyBF16Storage ? Builder.CreateBitCast(Arg, Ty) : Arg);
3314 Args.push_back(Builder.getInt32(
static_cast<int>(
RoundingMode)));
3316 Value *Rep = Builder.CreateIntrinsic(Ty, IID, Args);
3317 return IsLegacyBF16Storage ? Builder.CreateBitCast(Rep, CI->
getType()) : Rep;
3322 Value *Rep =
nullptr;
3324 if (Name ==
"abs.i" || Name ==
"abs.ll") {
3326 Rep = Builder.CreateIntrinsic(Intrinsic::abs, {Arg->
getType()},
3327 {Arg, Builder.getTrue()},
3329 }
else if (Name ==
"abs.bf16" || Name ==
"abs.bf16x2") {
3330 Type *Ty = (Name ==
"abs.bf16")
3334 Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg);
3335 Rep = Builder.CreateBitCast(Abs, CI->
getType());
3336 }
else if (Name ==
"fabs.f" || Name ==
"fabs.ftz.f" || Name ==
"fabs.d") {
3337 Intrinsic::ID IID = (Name ==
"fabs.ftz.f") ? Intrinsic::nvvm_fabs_ftz
3338 : Intrinsic::nvvm_fabs;
3339 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3340 }
else if (Name.consume_front(
"add.")) {
3343 }
else if (Name.consume_front(
"mul.")) {
3346 }
else if (Name.consume_front(
"ex2.approx.")) {
3348 Intrinsic::ID IID = Name.starts_with(
"ftz") ? Intrinsic::nvvm_ex2_approx_ftz
3349 : Intrinsic::nvvm_ex2_approx;
3350 Rep = Builder.CreateUnaryIntrinsic(IID, CI->
getArgOperand(0));
3351 }
else if (Name.consume_front(
"fma.")) {
3354 }
else if (Name.starts_with(
"atomic.load.add.f32.p") ||
3355 Name.starts_with(
"atomic.load.add.f64.p")) {
3358 Rep = Builder.CreateAtomicRMW(
3362 LLVMContext::MD_atomic_ignore_denormal_mode,
3367 }
else if (Name.starts_with(
"atomic.load.inc.32.p") ||
3368 Name.starts_with(
"atomic.load.dec.32.p")) {
3373 Rep = Builder.CreateAtomicRMW(
3377 }
else if (Name.starts_with(
"atomic.") && Name.contains(
".gen.")) {
3383 Op.contains(
".cta.") ?
"block" :
"");
3384 if (
Op.starts_with(
"cas.")) {
3386 Value *Pair = Builder.CreateAtomicCmpXchg(
3389 Rep = Builder.CreateExtractValue(Pair, 0);
3407 "unexpected nvvm scoped atomic intrinsic");
3408 Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val,
MaybeAlign(),
3412 LLVMContext::MD_atomic_ignore_denormal_mode,
3416 }
else if (Name ==
"clz.ll") {
3419 Value *Ctlz = Builder.CreateIntrinsic(Intrinsic::ctlz, {Arg->
getType()},
3420 {Arg, Builder.getFalse()},
3422 Rep = Builder.CreateTrunc(Ctlz, Builder.getInt32Ty(),
"ctlz.trunc");
3423 }
else if (Name ==
"popc.ll") {
3427 Value *Popc = Builder.CreateIntrinsic(Intrinsic::ctpop, {Arg->
getType()},
3428 Arg,
nullptr,
"ctpop");
3429 Rep = Builder.CreateTrunc(Popc, Builder.getInt32Ty(),
"ctpop.trunc");
3430 }
else if (Name ==
"h2f") {
3432 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
3433 Rep = Builder.CreateFPExt(Cast, Builder.getFloatTy());
3434 }
else if (Name.consume_front(
"bitcast.") &&
3435 (Name ==
"f2i" || Name ==
"i2f" || Name ==
"ll2d" ||
3438 }
else if (Name ==
"rotate.b32") {
3441 Rep = Builder.CreateIntrinsic(Builder.getInt32Ty(), Intrinsic::fshl,
3442 {Arg, Arg, ShiftAmt});
3443 }
else if (Name ==
"rotate.b64") {
3447 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3448 {Arg, Arg, ZExtShiftAmt});
3449 }
else if (Name ==
"rotate.right.b64") {
3453 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshr,
3454 {Arg, Arg, ZExtShiftAmt});
3455 }
else if (Name ==
"swap.lo.hi.b64") {
3458 Rep = Builder.CreateIntrinsic(Int64Ty, Intrinsic::fshl,
3459 {Arg, Arg, Builder.getInt64(32)});
3460 }
else if ((Name.consume_front(
"ptr.gen.to.") &&
3463 Name.starts_with(
".to.gen"))) {
3465 }
else if (Name.consume_front(
"ldg.global")) {
3469 Value *ASC = Builder.CreateAddrSpaceCast(Ptr, Builder.getPtrTy(1));
3472 LD->setMetadata(LLVMContext::MD_invariant_load, MD);
3474 }
else if (Name ==
"tanh.approx.f32") {
3478 Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->
getArgOperand(0),
3480 }
else if (Name ==
"barrier0" || Name ==
"barrier.n" || Name ==
"bar.sync") {
3482 Name.ends_with(
'0') ? Builder.getInt32(0) : CI->
getArgOperand(0);
3483 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
3485 }
else if (Name ==
"barrier") {
3486 Rep = Builder.CreateIntrinsic(
3487 Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
3489 }
else if (Name ==
"barrier.sync") {
3490 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
3492 }
else if (Name ==
"barrier.sync.cnt") {
3493 Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
3495 }
else if (Name ==
"barrier0.popc" || Name ==
"barrier0.and" ||
3496 Name ==
"barrier0.or") {
3498 C = Builder.CreateICmpNE(
C, Builder.getInt32(0));
3502 .
Case(
"barrier0.popc",
3503 Intrinsic::nvvm_barrier_cta_red_popc_aligned_all)
3504 .
Case(
"barrier0.and",
3505 Intrinsic::nvvm_barrier_cta_red_and_aligned_all)
3506 .
Case(
"barrier0.or",
3507 Intrinsic::nvvm_barrier_cta_red_or_aligned_all);
3508 Value *Bar = Builder.CreateIntrinsic(IID, {}, {Builder.getInt32(0),
C});
3509 Rep = Builder.CreateZExt(Bar, CI->
getType());
3523 ? Builder.CreateBitCast(Arg, NewType)
3526 Rep = Builder.CreateCall(NewFn, Args);
3527 if (
F->getReturnType()->isIntegerTy())
3528 Rep = Builder.CreateBitCast(Rep,
F->getReturnType());
3538 Value *Rep =
nullptr;
3540 if (Name.starts_with(
"sse4a.movnt.")) {
3552 Builder.CreateExtractElement(Arg1, (
uint64_t)0,
"extractelement");
3555 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3556 }
else if (Name.starts_with(
"avx.movnt.") ||
3557 Name.starts_with(
"avx512.storent.")) {
3569 SI->setMetadata(LLVMContext::MD_nontemporal,
Node);
3570 }
else if (Name ==
"sse2.storel.dq") {
3575 Value *BC0 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
3576 Value *Elt = Builder.CreateExtractElement(BC0, (
uint64_t)0);
3577 Builder.CreateAlignedStore(Elt, Arg0,
Align(1));
3578 }
else if (Name.starts_with(
"sse.storeu.") ||
3579 Name.starts_with(
"sse2.storeu.") ||
3580 Name.starts_with(
"avx.storeu.")) {
3583 Builder.CreateAlignedStore(Arg1, Arg0,
Align(1));
3584 }
else if (Name ==
"avx512.mask.store.ss") {
3588 }
else if (Name.starts_with(
"avx512.mask.store")) {
3590 bool Aligned = Name[17] !=
'u';
3593 }
else if (Name.starts_with(
"sse2.pcmp") || Name.starts_with(
"avx2.pcmp")) {
3596 bool CmpEq = Name[9] ==
'e';
3599 Rep = Builder.CreateSExt(Rep, CI->
getType(),
"");
3600 }
else if (Name.starts_with(
"avx512.broadcastm")) {
3607 Rep = Builder.CreateVectorSplat(NumElts, Rep);
3608 }
else if (Name ==
"sse.sqrt.ss" || Name ==
"sse2.sqrt.sd") {
3610 Value *Elt0 = Builder.CreateExtractElement(Vec, (
uint64_t)0);
3611 Elt0 = Builder.CreateIntrinsic(Intrinsic::sqrt, Elt0->
getType(), Elt0);
3612 Rep = Builder.CreateInsertElement(Vec, Elt0, (
uint64_t)0);
3613 }
else if (Name.starts_with(
"avx.sqrt.p") ||
3614 Name.starts_with(
"sse2.sqrt.p") ||
3615 Name.starts_with(
"sse.sqrt.p")) {
3616 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3617 {CI->getArgOperand(0)});
3618 }
else if (Name.starts_with(
"avx512.mask.sqrt.p")) {
3622 Intrinsic::ID IID = Name[18] ==
's' ? Intrinsic::x86_avx512_sqrt_ps_512
3623 : Intrinsic::x86_avx512_sqrt_pd_512;
3626 Rep = Builder.CreateIntrinsic(IID, Args);
3628 Rep = Builder.CreateIntrinsic(Intrinsic::sqrt, CI->
getType(),
3629 {CI->getArgOperand(0)});
3633 }
else if (Name.starts_with(
"avx512.ptestm") ||
3634 Name.starts_with(
"avx512.ptestnm")) {
3638 Rep = Builder.CreateAnd(Op0, Op1);
3644 Rep = Builder.CreateICmp(Pred, Rep, Zero);
3646 }
else if (Name.starts_with(
"avx512.mask.pbroadcast")) {
3649 Rep = Builder.CreateVectorSplat(NumElts, CI->
getArgOperand(0));
3652 }
else if (Name.starts_with(
"avx512.kunpck")) {
3657 for (
unsigned i = 0; i != NumElts; ++i)
3666 Rep = Builder.CreateShuffleVector(
RHS,
LHS,
ArrayRef(Indices, NumElts));
3667 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3668 }
else if (Name ==
"avx512.kand.w") {
3671 Rep = Builder.CreateAnd(
LHS,
RHS);
3672 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3673 }
else if (Name ==
"avx512.kandn.w") {
3676 LHS = Builder.CreateNot(
LHS);
3677 Rep = Builder.CreateAnd(
LHS,
RHS);
3678 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3679 }
else if (Name ==
"avx512.kor.w") {
3682 Rep = Builder.CreateOr(
LHS,
RHS);
3683 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3684 }
else if (Name ==
"avx512.kxor.w") {
3687 Rep = Builder.CreateXor(
LHS,
RHS);
3688 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3689 }
else if (Name ==
"avx512.kxnor.w") {
3692 LHS = Builder.CreateNot(
LHS);
3693 Rep = Builder.CreateXor(
LHS,
RHS);
3694 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3695 }
else if (Name ==
"avx512.knot.w") {
3697 Rep = Builder.CreateNot(Rep);
3698 Rep = Builder.CreateBitCast(Rep, CI->
getType());
3699 }
else if (Name ==
"avx512.kortestz.w" || Name ==
"avx512.kortestc.w") {
3702 Rep = Builder.CreateOr(
LHS,
RHS);
3703 Rep = Builder.CreateBitCast(Rep, Builder.getInt16Ty());
3705 if (Name[14] ==
'c')
3709 Rep = Builder.CreateICmpEQ(Rep,
C);
3710 Rep = Builder.CreateZExt(Rep, Builder.getInt32Ty());
3711 }
else if (Name ==
"sse.add.ss" || Name ==
"sse2.add.sd" ||
3712 Name ==
"sse.sub.ss" || Name ==
"sse2.sub.sd" ||
3713 Name ==
"sse.mul.ss" || Name ==
"sse2.mul.sd" ||
3714 Name ==
"sse.div.ss" || Name ==
"sse2.div.sd") {
3717 ConstantInt::get(I32Ty, 0));
3719 ConstantInt::get(I32Ty, 0));
3721 if (Name.contains(
".add."))
3722 EltOp = Builder.CreateFAdd(Elt0, Elt1);
3723 else if (Name.contains(
".sub."))
3724 EltOp = Builder.CreateFSub(Elt0, Elt1);
3725 else if (Name.contains(
".mul."))
3726 EltOp = Builder.CreateFMul(Elt0, Elt1);
3728 EltOp = Builder.CreateFDiv(Elt0, Elt1);
3729 Rep = Builder.CreateInsertElement(CI->
getArgOperand(0), EltOp,
3730 ConstantInt::get(I32Ty, 0));
3731 }
else if (Name.starts_with(
"avx512.mask.pcmp")) {
3733 bool CmpEq = Name[16] ==
'e';
3735 }
else if (Name.starts_with(
"avx512.mask.vpshufbitqmb.")) {
3737 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3744 IID = Intrinsic::x86_avx512_vpshufbitqmb_128;
3747 IID = Intrinsic::x86_avx512_vpshufbitqmb_256;
3750 IID = Intrinsic::x86_avx512_vpshufbitqmb_512;
3757 }
else if (Name.starts_with(
"avx512.mask.fpclass.p")) {
3759 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3760 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3762 if (VecWidth == 128 && EltWidth == 32)
3763 IID = Intrinsic::x86_avx512_fpclass_ps_128;
3764 else if (VecWidth == 256 && EltWidth == 32)
3765 IID = Intrinsic::x86_avx512_fpclass_ps_256;
3766 else if (VecWidth == 512 && EltWidth == 32)
3767 IID = Intrinsic::x86_avx512_fpclass_ps_512;
3768 else if (VecWidth == 128 && EltWidth == 64)
3769 IID = Intrinsic::x86_avx512_fpclass_pd_128;
3770 else if (VecWidth == 256 && EltWidth == 64)
3771 IID = Intrinsic::x86_avx512_fpclass_pd_256;
3772 else if (VecWidth == 512 && EltWidth == 64)
3773 IID = Intrinsic::x86_avx512_fpclass_pd_512;
3780 }
else if (Name.starts_with(
"avx512.cmp.p")) {
3783 unsigned VecWidth =
OpTy->getPrimitiveSizeInBits();
3784 unsigned EltWidth =
OpTy->getScalarSizeInBits();
3786 if (VecWidth == 128 && EltWidth == 32)
3787 IID = Intrinsic::x86_avx512_mask_cmp_ps_128;
3788 else if (VecWidth == 256 && EltWidth == 32)
3789 IID = Intrinsic::x86_avx512_mask_cmp_ps_256;
3790 else if (VecWidth == 512 && EltWidth == 32)
3791 IID = Intrinsic::x86_avx512_mask_cmp_ps_512;
3792 else if (VecWidth == 128 && EltWidth == 64)
3793 IID = Intrinsic::x86_avx512_mask_cmp_pd_128;
3794 else if (VecWidth == 256 && EltWidth == 64)
3795 IID = Intrinsic::x86_avx512_mask_cmp_pd_256;
3796 else if (VecWidth == 512 && EltWidth == 64)
3797 IID = Intrinsic::x86_avx512_mask_cmp_pd_512;
3802 if (VecWidth == 512)
3804 Args.push_back(Mask);
3806 Rep = Builder.CreateIntrinsic(IID, Args);
3807 }
else if (Name.starts_with(
"avx512.mask.cmp.")) {
3811 }
else if (Name.starts_with(
"avx512.mask.ucmp.")) {
3814 }
else if (Name.starts_with(
"avx512.cvtb2mask.") ||
3815 Name.starts_with(
"avx512.cvtw2mask.") ||
3816 Name.starts_with(
"avx512.cvtd2mask.") ||
3817 Name.starts_with(
"avx512.cvtq2mask.")) {
3822 }
else if (Name ==
"ssse3.pabs.b.128" || Name ==
"ssse3.pabs.w.128" ||
3823 Name ==
"ssse3.pabs.d.128" || Name.starts_with(
"avx2.pabs") ||
3824 Name.starts_with(
"avx512.mask.pabs")) {
3826 }
else if (Name ==
"sse41.pmaxsb" || Name ==
"sse2.pmaxs.w" ||
3827 Name ==
"sse41.pmaxsd" || Name.starts_with(
"avx2.pmaxs") ||
3828 Name.starts_with(
"avx512.mask.pmaxs")) {
3830 }
else if (Name ==
"sse2.pmaxu.b" || Name ==
"sse41.pmaxuw" ||
3831 Name ==
"sse41.pmaxud" || Name.starts_with(
"avx2.pmaxu") ||
3832 Name.starts_with(
"avx512.mask.pmaxu")) {
3834 }
else if (Name ==
"sse41.pminsb" || Name ==
"sse2.pmins.w" ||
3835 Name ==
"sse41.pminsd" || Name.starts_with(
"avx2.pmins") ||
3836 Name.starts_with(
"avx512.mask.pmins")) {
3838 }
else if (Name ==
"sse2.pminu.b" || Name ==
"sse41.pminuw" ||
3839 Name ==
"sse41.pminud" || Name.starts_with(
"avx2.pminu") ||
3840 Name.starts_with(
"avx512.mask.pminu")) {
3842 }
else if (Name ==
"sse2.pmulh.w" || Name.starts_with(
"avx2.pmulh.w") ||
3843 Name.starts_with(
"avx512.pmulh.w")) {
3845 }
else if (Name ==
"sse2.pmulhu.w" || Name.starts_with(
"avx2.pmulhu.w") ||
3846 Name.starts_with(
"avx512.pmulhu.w")) {
3848 }
else if (Name ==
"sse2.pmulu.dq" || Name ==
"avx2.pmulu.dq" ||
3849 Name ==
"avx512.pmulu.dq.512" ||
3850 Name.starts_with(
"avx512.mask.pmulu.dq.")) {
3852 }
else if (Name ==
"sse41.pmuldq" || Name ==
"avx2.pmul.dq" ||
3853 Name ==
"avx512.pmul.dq.512" ||
3854 Name.starts_with(
"avx512.mask.pmul.dq.")) {
3856 }
else if (Name ==
"sse.cvtsi2ss" || Name ==
"sse2.cvtsi2sd" ||
3857 Name ==
"sse.cvtsi642ss" || Name ==
"sse2.cvtsi642sd") {
3862 }
else if (Name ==
"avx512.cvtusi2sd") {
3867 }
else if (Name ==
"sse2.cvtss2sd") {
3869 Rep = Builder.CreateFPExt(
3872 }
else if (Name ==
"sse2.cvtdq2pd" || Name ==
"sse2.cvtdq2ps" ||
3873 Name ==
"avx.cvtdq2.pd.256" || Name ==
"avx.cvtdq2.ps.256" ||
3874 Name.starts_with(
"avx512.mask.cvtdq2pd.") ||
3875 Name.starts_with(
"avx512.mask.cvtudq2pd.") ||
3876 Name.starts_with(
"avx512.mask.cvtdq2ps.") ||
3877 Name.starts_with(
"avx512.mask.cvtudq2ps.") ||
3878 Name.starts_with(
"avx512.mask.cvtqq2pd.") ||
3879 Name.starts_with(
"avx512.mask.cvtuqq2pd.") ||
3880 Name ==
"avx512.mask.cvtqq2ps.256" ||
3881 Name ==
"avx512.mask.cvtqq2ps.512" ||
3882 Name ==
"avx512.mask.cvtuqq2ps.256" ||
3883 Name ==
"avx512.mask.cvtuqq2ps.512" || Name ==
"sse2.cvtps2pd" ||
3884 Name ==
"avx.cvt.ps2.pd.256" ||
3885 Name ==
"avx512.mask.cvtps2pd.128" ||
3886 Name ==
"avx512.mask.cvtps2pd.256") {
3891 unsigned NumDstElts = DstTy->getNumElements();
3892 if (NumDstElts < SrcTy->getNumElements()) {
3893 assert(NumDstElts == 2 &&
"Unexpected vector size");
3894 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1});
3897 bool IsPS2PD = SrcTy->getElementType()->isFloatTy();
3898 bool IsUnsigned = Name.contains(
"cvtu");
3900 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtps2pd");
3904 Intrinsic::ID IID = IsUnsigned ? Intrinsic::x86_avx512_uitofp_round
3905 : Intrinsic::x86_avx512_sitofp_round;
3906 Rep = Builder.CreateIntrinsic(IID, {DstTy, SrcTy},
3909 Rep = IsUnsigned ? Builder.CreateUIToFP(Rep, DstTy,
"cvt")
3910 : Builder.CreateSIToFP(Rep, DstTy,
"cvt");
3916 }
else if (Name.starts_with(
"avx512.mask.vcvtph2ps.") ||
3917 Name.starts_with(
"vcvtph2ps.")) {
3921 unsigned NumDstElts = DstTy->getNumElements();
3922 if (NumDstElts != SrcTy->getNumElements()) {
3923 assert(NumDstElts == 4 &&
"Unexpected vector size");
3924 Rep = Builder.CreateShuffleVector(Rep, Rep,
ArrayRef<int>{0, 1, 2, 3});
3926 Rep = Builder.CreateBitCast(
3928 Rep = Builder.CreateFPExt(Rep, DstTy,
"cvtph2ps");
3932 }
else if (Name.starts_with(
"avx512.mask.load")) {
3934 bool Aligned = Name[16] !=
'u';
3937 }
else if (Name.starts_with(
"avx512.mask.expand.load.")) {
3941 ResultTy->getNumElements());
3942 Rep = Builder.CreateIntrinsic(
3943 Intrinsic::masked_expandload, {ResultTy, PtrTy},
3945 }
else if (Name.starts_with(
"avx512.mask.compress.store.")) {
3951 Rep = Builder.CreateIntrinsic(
3952 Intrinsic::masked_compressstore, {ResultTy, PtrTy},
3954 }
else if (Name.starts_with(
"avx512.mask.compress.") ||
3955 Name.starts_with(
"avx512.mask.expand.")) {
3959 ResultTy->getNumElements());
3961 bool IsCompress = Name[12] ==
'c';
3962 Intrinsic::ID IID = IsCompress ? Intrinsic::x86_avx512_mask_compress
3963 : Intrinsic::x86_avx512_mask_expand;
3964 Rep = Builder.CreateIntrinsic(
3966 }
else if (Name.starts_with(
"xop.vpcom")) {
3968 if (Name.ends_with(
"ub") || Name.ends_with(
"uw") || Name.ends_with(
"ud") ||
3969 Name.ends_with(
"uq"))
3971 else if (Name.ends_with(
"b") || Name.ends_with(
"w") ||
3972 Name.ends_with(
"d") || Name.ends_with(
"q"))
3981 Name = Name.substr(9);
3982 if (Name.starts_with(
"lt"))
3984 else if (Name.starts_with(
"le"))
3986 else if (Name.starts_with(
"gt"))
3988 else if (Name.starts_with(
"ge"))
3990 else if (Name.starts_with(
"eq"))
3992 else if (Name.starts_with(
"ne"))
3994 else if (Name.starts_with(
"false"))
3996 else if (Name.starts_with(
"true"))
4003 }
else if (Name.starts_with(
"xop.vpcmov")) {
4005 Value *NotSel = Builder.CreateNot(Sel);
4008 Rep = Builder.CreateOr(Sel0, Sel1);
4009 }
else if (Name.starts_with(
"xop.vprot") || Name.starts_with(
"avx512.prol") ||
4010 Name.starts_with(
"avx512.mask.prol")) {
4012 }
else if (Name.starts_with(
"avx512.pror") ||
4013 Name.starts_with(
"avx512.mask.pror")) {
4015 }
else if (Name.starts_with(
"avx512.vpshld.") ||
4016 Name.starts_with(
"avx512.mask.vpshld") ||
4017 Name.starts_with(
"avx512.maskz.vpshld")) {
4018 bool ZeroMask = Name[11] ==
'z';
4020 }
else if (Name.starts_with(
"avx512.vpshrd.") ||
4021 Name.starts_with(
"avx512.mask.vpshrd") ||
4022 Name.starts_with(
"avx512.maskz.vpshrd")) {
4023 bool ZeroMask = Name[11] ==
'z';
4025 }
else if (Name ==
"sse42.crc32.64.8") {
4028 Rep = Builder.CreateIntrinsic(Intrinsic::x86_sse42_crc32_32_8,
4030 Rep = Builder.CreateZExt(Rep, CI->
getType(),
"");
4031 }
else if (Name.starts_with(
"avx.vbroadcast.s") ||
4032 Name.starts_with(
"avx512.vbroadcast.s")) {
4035 Type *EltTy = VecTy->getElementType();
4036 unsigned EltNum = VecTy->getNumElements();
4040 for (
unsigned I = 0;
I < EltNum; ++
I)
4041 Rep = Builder.CreateInsertElement(Rep,
Load, ConstantInt::get(I32Ty,
I));
4042 }
else if (Name.starts_with(
"sse41.pmovsx") ||
4043 Name.starts_with(
"sse41.pmovzx") ||
4044 Name.starts_with(
"avx2.pmovsx") ||
4045 Name.starts_with(
"avx2.pmovzx") ||
4046 Name.starts_with(
"avx512.mask.pmovsx") ||
4047 Name.starts_with(
"avx512.mask.pmovzx")) {
4049 unsigned NumDstElts = DstTy->getNumElements();
4053 for (
unsigned i = 0; i != NumDstElts; ++i)
4058 bool DoSext = Name.contains(
"pmovsx");
4060 DoSext ? Builder.CreateSExt(SV, DstTy) : Builder.CreateZExt(SV, DstTy);
4065 }
else if (Name ==
"avx512.mask.pmov.qd.256" ||
4066 Name ==
"avx512.mask.pmov.qd.512" ||
4067 Name ==
"avx512.mask.pmov.wb.256" ||
4068 Name ==
"avx512.mask.pmov.wb.512") {
4073 }
else if (Name.starts_with(
"avx.vbroadcastf128") ||
4074 Name ==
"avx2.vbroadcasti128") {
4080 if (NumSrcElts == 2)
4083 Rep = Builder.CreateShuffleVector(
Load,
4085 }
else if (Name.starts_with(
"avx512.mask.shuf.i") ||
4086 Name.starts_with(
"avx512.mask.shuf.f")) {
4091 unsigned ControlBitsMask = NumLanes - 1;
4092 unsigned NumControlBits = NumLanes / 2;
4095 for (
unsigned l = 0; l != NumLanes; ++l) {
4096 unsigned LaneMask = (
Imm >> (l * NumControlBits)) & ControlBitsMask;
4098 if (l >= NumLanes / 2)
4099 LaneMask += NumLanes;
4100 for (
unsigned i = 0; i != NumElementsInLane; ++i)
4101 ShuffleMask.push_back(LaneMask * NumElementsInLane + i);
4107 }
else if (Name.starts_with(
"avx512.mask.broadcastf") ||
4108 Name.starts_with(
"avx512.mask.broadcasti")) {
4111 unsigned NumDstElts =
4115 for (
unsigned i = 0; i != NumDstElts; ++i)
4116 ShuffleMask[i] = i % NumSrcElts;
4122 }
else if (Name.starts_with(
"avx2.pbroadcast") ||
4123 Name.starts_with(
"avx2.vbroadcast") ||
4124 Name.starts_with(
"avx512.pbroadcast") ||
4125 Name.starts_with(
"avx512.mask.broadcast.s")) {
4132 Rep = Builder.CreateShuffleVector(
Op, M);
4137 }
else if (Name.starts_with(
"sse2.padds.") ||
4138 Name.starts_with(
"avx2.padds.") ||
4139 Name.starts_with(
"avx512.padds.") ||
4140 Name.starts_with(
"avx512.mask.padds.")) {
4142 }
else if (Name.starts_with(
"sse2.psubs.") ||
4143 Name.starts_with(
"avx2.psubs.") ||
4144 Name.starts_with(
"avx512.psubs.") ||
4145 Name.starts_with(
"avx512.mask.psubs.")) {
4147 }
else if (Name.starts_with(
"sse2.paddus.") ||
4148 Name.starts_with(
"avx2.paddus.") ||
4149 Name.starts_with(
"avx512.mask.paddus.")) {
4151 }
else if (Name.starts_with(
"sse2.psubus.") ||
4152 Name.starts_with(
"avx2.psubus.") ||
4153 Name.starts_with(
"avx512.mask.psubus.")) {
4155 }
else if (Name.starts_with(
"avx512.mask.palignr.")) {
4160 }
else if (Name.starts_with(
"avx512.mask.valign.")) {
4164 }
else if (Name ==
"sse2.psll.dq" || Name ==
"avx2.psll.dq") {
4169 }
else if (Name ==
"sse2.psrl.dq" || Name ==
"avx2.psrl.dq") {
4174 }
else if (Name ==
"sse2.psll.dq.bs" || Name ==
"avx2.psll.dq.bs" ||
4175 Name ==
"avx512.psll.dq.512") {
4179 }
else if (Name ==
"sse2.psrl.dq.bs" || Name ==
"avx2.psrl.dq.bs" ||
4180 Name ==
"avx512.psrl.dq.512") {
4184 }
else if (Name ==
"sse41.pblendw" || Name.starts_with(
"sse41.blendp") ||
4185 Name.starts_with(
"avx.blend.p") || Name ==
"avx2.pblendw" ||
4186 Name.starts_with(
"avx2.pblendd.")) {
4191 unsigned NumElts = VecTy->getNumElements();
4194 for (
unsigned i = 0; i != NumElts; ++i)
4195 Idxs[i] = ((
Imm >> (i % 8)) & 1) ? i + NumElts : i;
4197 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4198 }
else if (Name.starts_with(
"avx.vinsertf128.") ||
4199 Name ==
"avx2.vinserti128" ||
4200 Name.starts_with(
"avx512.mask.insert")) {
4204 unsigned DstNumElts =
4206 unsigned SrcNumElts =
4208 unsigned Scale = DstNumElts / SrcNumElts;
4215 for (
unsigned i = 0; i != SrcNumElts; ++i)
4217 for (
unsigned i = SrcNumElts; i != DstNumElts; ++i)
4218 Idxs[i] = SrcNumElts;
4219 Rep = Builder.CreateShuffleVector(Op1, Idxs);
4233 for (
unsigned i = 0; i != DstNumElts; ++i)
4236 for (
unsigned i = 0; i != SrcNumElts; ++i)
4237 Idxs[i +
Imm * SrcNumElts] = i + DstNumElts;
4238 Rep = Builder.CreateShuffleVector(Op0, Rep, Idxs);
4244 }
else if (Name.starts_with(
"avx.vextractf128.") ||
4245 Name ==
"avx2.vextracti128" ||
4246 Name.starts_with(
"avx512.mask.vextract")) {
4249 unsigned DstNumElts =
4251 unsigned SrcNumElts =
4253 unsigned Scale = SrcNumElts / DstNumElts;
4260 for (
unsigned i = 0; i != DstNumElts; ++i) {
4261 Idxs[i] = i + (
Imm * DstNumElts);
4263 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4269 }
else if (Name.starts_with(
"avx512.mask.perm.df.") ||
4270 Name.starts_with(
"avx512.mask.perm.di.")) {
4274 unsigned NumElts = VecTy->getNumElements();
4277 for (
unsigned i = 0; i != NumElts; ++i)
4278 Idxs[i] = (i & ~0x3) + ((
Imm >> (2 * (i & 0x3))) & 3);
4280 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4285 }
else if (Name.starts_with(
"avx.vperm2f128.") || Name ==
"avx2.vperm2i128") {
4297 unsigned HalfSize = NumElts / 2;
4309 unsigned StartIndex = (
Imm & 0x01) ? HalfSize : 0;
4310 for (
unsigned i = 0; i < HalfSize; ++i)
4311 ShuffleMask[i] = StartIndex + i;
4314 StartIndex = (
Imm & 0x10) ? HalfSize : 0;
4315 for (
unsigned i = 0; i < HalfSize; ++i)
4316 ShuffleMask[i + HalfSize] = NumElts + StartIndex + i;
4318 Rep = Builder.CreateShuffleVector(V0,
V1, ShuffleMask);
4320 }
else if (Name.starts_with(
"avx.vpermil.") || Name ==
"sse2.pshuf.d" ||
4321 Name.starts_with(
"avx512.mask.vpermil.p") ||
4322 Name.starts_with(
"avx512.mask.pshuf.d.")) {
4326 unsigned NumElts = VecTy->getNumElements();
4328 unsigned IdxSize = 64 / VecTy->getScalarSizeInBits();
4329 unsigned IdxMask = ((1 << IdxSize) - 1);
4335 for (
unsigned i = 0; i != NumElts; ++i)
4336 Idxs[i] = ((
Imm >> ((i * IdxSize) % 8)) & IdxMask) | (i & ~IdxMask);
4338 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4343 }
else if (Name ==
"sse2.pshufl.w" ||
4344 Name.starts_with(
"avx512.mask.pshufl.w.")) {
4349 if (Name ==
"sse2.pshufl.w" && NumElts % 8 != 0)
4353 for (
unsigned l = 0; l != NumElts; l += 8) {
4354 for (
unsigned i = 0; i != 4; ++i)
4355 Idxs[i + l] = ((
Imm >> (2 * i)) & 0x3) + l;
4356 for (
unsigned i = 4; i != 8; ++i)
4357 Idxs[i + l] = i + l;
4360 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4365 }
else if (Name ==
"sse2.pshufh.w" ||
4366 Name.starts_with(
"avx512.mask.pshufh.w.")) {
4371 if (Name ==
"sse2.pshufh.w" && NumElts % 8 != 0)
4375 for (
unsigned l = 0; l != NumElts; l += 8) {
4376 for (
unsigned i = 0; i != 4; ++i)
4377 Idxs[i + l] = i + l;
4378 for (
unsigned i = 0; i != 4; ++i)
4379 Idxs[i + l + 4] = ((
Imm >> (2 * i)) & 0x3) + 4 + l;
4382 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4387 }
else if (Name.starts_with(
"avx512.mask.shuf.p")) {
4394 unsigned HalfLaneElts = NumLaneElts / 2;
4397 for (
unsigned i = 0; i != NumElts; ++i) {
4399 Idxs[i] = i - (i % NumLaneElts);
4401 if ((i % NumLaneElts) >= HalfLaneElts)
4405 Idxs[i] += (
Imm >> ((i * HalfLaneElts) % 8)) & ((1 << HalfLaneElts) - 1);
4408 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4412 }
else if (Name.starts_with(
"avx512.mask.movddup") ||
4413 Name.starts_with(
"avx512.mask.movshdup") ||
4414 Name.starts_with(
"avx512.mask.movsldup")) {
4420 if (Name.starts_with(
"avx512.mask.movshdup."))
4424 for (
unsigned l = 0; l != NumElts; l += NumLaneElts)
4425 for (
unsigned i = 0; i != NumLaneElts; i += 2) {
4426 Idxs[i + l + 0] = i + l +
Offset;
4427 Idxs[i + l + 1] = i + l +
Offset;
4430 Rep = Builder.CreateShuffleVector(Op0, Op0, Idxs);
4434 }
else if (Name.starts_with(
"avx512.mask.punpckl") ||
4435 Name.starts_with(
"avx512.mask.unpckl.")) {
4442 for (
int l = 0; l != NumElts; l += NumLaneElts)
4443 for (
int i = 0; i != NumLaneElts; ++i)
4444 Idxs[i + l] = l + (i / 2) + NumElts * (i % 2);
4446 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4450 }
else if (Name.starts_with(
"avx512.mask.punpckh") ||
4451 Name.starts_with(
"avx512.mask.unpckh.")) {
4458 for (
int l = 0; l != NumElts; l += NumLaneElts)
4459 for (
int i = 0; i != NumLaneElts; ++i)
4460 Idxs[i + l] = (NumLaneElts / 2) + l + (i / 2) + NumElts * (i % 2);
4462 Rep = Builder.CreateShuffleVector(Op0, Op1, Idxs);
4466 }
else if (Name.starts_with(
"avx512.mask.and.") ||
4467 Name.starts_with(
"avx512.mask.pand.")) {
4470 Rep = Builder.CreateAnd(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4472 Rep = Builder.CreateBitCast(Rep, FTy);
4475 }
else if (Name.starts_with(
"avx512.mask.andn.") ||
4476 Name.starts_with(
"avx512.mask.pandn.")) {
4479 Rep = Builder.CreateNot(Builder.CreateBitCast(CI->
getArgOperand(0), ITy));
4480 Rep = Builder.CreateAnd(Rep,
4482 Rep = Builder.CreateBitCast(Rep, FTy);
4485 }
else if (Name.starts_with(
"avx512.mask.or.") ||
4486 Name.starts_with(
"avx512.mask.por.")) {
4489 Rep = Builder.CreateOr(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4491 Rep = Builder.CreateBitCast(Rep, FTy);
4494 }
else if (Name.starts_with(
"avx512.mask.xor.") ||
4495 Name.starts_with(
"avx512.mask.pxor.")) {
4498 Rep = Builder.CreateXor(Builder.CreateBitCast(CI->
getArgOperand(0), ITy),
4500 Rep = Builder.CreateBitCast(Rep, FTy);
4503 }
else if (Name.starts_with(
"avx512.mask.padd.")) {
4507 }
else if (Name.starts_with(
"avx512.mask.psub.")) {
4511 }
else if (Name.starts_with(
"avx512.mask.pmull.")) {
4515 }
else if (Name.starts_with(
"avx512.mask.add.p")) {
4516 if (Name.ends_with(
".512")) {
4518 if (Name[17] ==
's')
4519 IID = Intrinsic::x86_avx512_add_ps_512;
4521 IID = Intrinsic::x86_avx512_add_pd_512;
4523 Rep = Builder.CreateIntrinsic(
4531 }
else if (Name.starts_with(
"avx512.mask.div.p")) {
4532 if (Name.ends_with(
".512")) {
4534 if (Name[17] ==
's')
4535 IID = Intrinsic::x86_avx512_div_ps_512;
4537 IID = Intrinsic::x86_avx512_div_pd_512;
4539 Rep = Builder.CreateIntrinsic(
4547 }
else if (Name.starts_with(
"avx512.mask.mul.p")) {
4548 if (Name.ends_with(
".512")) {
4550 if (Name[17] ==
's')
4551 IID = Intrinsic::x86_avx512_mul_ps_512;
4553 IID = Intrinsic::x86_avx512_mul_pd_512;
4555 Rep = Builder.CreateIntrinsic(
4563 }
else if (Name.starts_with(
"avx512.mask.sub.p")) {
4564 if (Name.ends_with(
".512")) {
4566 if (Name[17] ==
's')
4567 IID = Intrinsic::x86_avx512_sub_ps_512;
4569 IID = Intrinsic::x86_avx512_sub_pd_512;
4571 Rep = Builder.CreateIntrinsic(
4579 }
else if ((Name.starts_with(
"avx512.mask.max.p") ||
4580 Name.starts_with(
"avx512.mask.min.p")) &&
4581 Name.drop_front(18) ==
".512") {
4582 bool IsDouble = Name[17] ==
'd';
4583 bool IsMin = Name[13] ==
'i';
4585 {Intrinsic::x86_avx512_max_ps_512, Intrinsic::x86_avx512_max_pd_512},
4586 {Intrinsic::x86_avx512_min_ps_512, Intrinsic::x86_avx512_min_pd_512}};
4589 Rep = Builder.CreateIntrinsic(
4594 }
else if (Name.starts_with(
"avx512.mask.lzcnt.")) {
4596 Builder.CreateIntrinsic(Intrinsic::ctlz, CI->
getType(),
4597 {CI->getArgOperand(0), Builder.getInt1(false)});
4600 }
else if (Name.starts_with(
"avx512.mask.psll")) {
4601 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4602 bool IsVariable = Name[16] ==
'v';
4603 char Size = Name[16] ==
'.' ? Name[17]
4604 : Name[17] ==
'.' ? Name[18]
4605 : Name[18] ==
'.' ? Name[19]
4609 if (IsVariable && Name[17] !=
'.') {
4610 if (
Size ==
'd' && Name[17] ==
'2')
4611 IID = Intrinsic::x86_avx2_psllv_q;
4612 else if (
Size ==
'd' && Name[17] ==
'4')
4613 IID = Intrinsic::x86_avx2_psllv_q_256;
4614 else if (
Size ==
's' && Name[17] ==
'4')
4615 IID = Intrinsic::x86_avx2_psllv_d;
4616 else if (
Size ==
's' && Name[17] ==
'8')
4617 IID = Intrinsic::x86_avx2_psllv_d_256;
4618 else if (
Size ==
'h' && Name[17] ==
'8')
4619 IID = Intrinsic::x86_avx512_psllv_w_128;
4620 else if (
Size ==
'h' && Name[17] ==
'1')
4621 IID = Intrinsic::x86_avx512_psllv_w_256;
4622 else if (Name[17] ==
'3' && Name[18] ==
'2')
4623 IID = Intrinsic::x86_avx512_psllv_w_512;
4626 }
else if (Name.ends_with(
".128")) {
4628 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_d
4629 : Intrinsic::x86_sse2_psll_d;
4630 else if (
Size ==
'q')
4631 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_q
4632 : Intrinsic::x86_sse2_psll_q;
4633 else if (
Size ==
'w')
4634 IID = IsImmediate ? Intrinsic::x86_sse2_pslli_w
4635 : Intrinsic::x86_sse2_psll_w;
4638 }
else if (Name.ends_with(
".256")) {
4640 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_d
4641 : Intrinsic::x86_avx2_psll_d;
4642 else if (
Size ==
'q')
4643 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_q
4644 : Intrinsic::x86_avx2_psll_q;
4645 else if (
Size ==
'w')
4646 IID = IsImmediate ? Intrinsic::x86_avx2_pslli_w
4647 : Intrinsic::x86_avx2_psll_w;
4652 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_d_512
4653 : IsVariable ? Intrinsic::x86_avx512_psllv_d_512
4654 : Intrinsic::x86_avx512_psll_d_512;
4655 else if (
Size ==
'q')
4656 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_q_512
4657 : IsVariable ? Intrinsic::x86_avx512_psllv_q_512
4658 : Intrinsic::x86_avx512_psll_q_512;
4659 else if (
Size ==
'w')
4660 IID = IsImmediate ? Intrinsic::x86_avx512_pslli_w_512
4661 : Intrinsic::x86_avx512_psll_w_512;
4667 }
else if (Name.starts_with(
"avx512.mask.psrl")) {
4668 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4669 bool IsVariable = Name[16] ==
'v';
4670 char Size = Name[16] ==
'.' ? Name[17]
4671 : Name[17] ==
'.' ? Name[18]
4672 : Name[18] ==
'.' ? Name[19]
4676 if (IsVariable && Name[17] !=
'.') {
4677 if (
Size ==
'd' && Name[17] ==
'2')
4678 IID = Intrinsic::x86_avx2_psrlv_q;
4679 else if (
Size ==
'd' && Name[17] ==
'4')
4680 IID = Intrinsic::x86_avx2_psrlv_q_256;
4681 else if (
Size ==
's' && Name[17] ==
'4')
4682 IID = Intrinsic::x86_avx2_psrlv_d;
4683 else if (
Size ==
's' && Name[17] ==
'8')
4684 IID = Intrinsic::x86_avx2_psrlv_d_256;
4685 else if (
Size ==
'h' && Name[17] ==
'8')
4686 IID = Intrinsic::x86_avx512_psrlv_w_128;
4687 else if (
Size ==
'h' && Name[17] ==
'1')
4688 IID = Intrinsic::x86_avx512_psrlv_w_256;
4689 else if (Name[17] ==
'3' && Name[18] ==
'2')
4690 IID = Intrinsic::x86_avx512_psrlv_w_512;
4693 }
else if (Name.ends_with(
".128")) {
4695 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_d
4696 : Intrinsic::x86_sse2_psrl_d;
4697 else if (
Size ==
'q')
4698 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_q
4699 : Intrinsic::x86_sse2_psrl_q;
4700 else if (
Size ==
'w')
4701 IID = IsImmediate ? Intrinsic::x86_sse2_psrli_w
4702 : Intrinsic::x86_sse2_psrl_w;
4705 }
else if (Name.ends_with(
".256")) {
4707 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_d
4708 : Intrinsic::x86_avx2_psrl_d;
4709 else if (
Size ==
'q')
4710 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_q
4711 : Intrinsic::x86_avx2_psrl_q;
4712 else if (
Size ==
'w')
4713 IID = IsImmediate ? Intrinsic::x86_avx2_psrli_w
4714 : Intrinsic::x86_avx2_psrl_w;
4719 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_d_512
4720 : IsVariable ? Intrinsic::x86_avx512_psrlv_d_512
4721 : Intrinsic::x86_avx512_psrl_d_512;
4722 else if (
Size ==
'q')
4723 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_q_512
4724 : IsVariable ? Intrinsic::x86_avx512_psrlv_q_512
4725 : Intrinsic::x86_avx512_psrl_q_512;
4726 else if (
Size ==
'w')
4727 IID = IsImmediate ? Intrinsic::x86_avx512_psrli_w_512
4728 : Intrinsic::x86_avx512_psrl_w_512;
4734 }
else if (Name.starts_with(
"avx512.mask.psra")) {
4735 bool IsImmediate = Name[16] ==
'i' || (Name.size() > 18 && Name[18] ==
'i');
4736 bool IsVariable = Name[16] ==
'v';
4737 char Size = Name[16] ==
'.' ? Name[17]
4738 : Name[17] ==
'.' ? Name[18]
4739 : Name[18] ==
'.' ? Name[19]
4743 if (IsVariable && Name[17] !=
'.') {
4744 if (
Size ==
's' && Name[17] ==
'4')
4745 IID = Intrinsic::x86_avx2_psrav_d;
4746 else if (
Size ==
's' && Name[17] ==
'8')
4747 IID = Intrinsic::x86_avx2_psrav_d_256;
4748 else if (
Size ==
'h' && Name[17] ==
'8')
4749 IID = Intrinsic::x86_avx512_psrav_w_128;
4750 else if (
Size ==
'h' && Name[17] ==
'1')
4751 IID = Intrinsic::x86_avx512_psrav_w_256;
4752 else if (Name[17] ==
'3' && Name[18] ==
'2')
4753 IID = Intrinsic::x86_avx512_psrav_w_512;
4756 }
else if (Name.ends_with(
".128")) {
4758 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_d
4759 : Intrinsic::x86_sse2_psra_d;
4760 else if (
Size ==
'q')
4761 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_128
4762 : IsVariable ? Intrinsic::x86_avx512_psrav_q_128
4763 : Intrinsic::x86_avx512_psra_q_128;
4764 else if (
Size ==
'w')
4765 IID = IsImmediate ? Intrinsic::x86_sse2_psrai_w
4766 : Intrinsic::x86_sse2_psra_w;
4769 }
else if (Name.ends_with(
".256")) {
4771 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_d
4772 : Intrinsic::x86_avx2_psra_d;
4773 else if (
Size ==
'q')
4774 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_256
4775 : IsVariable ? Intrinsic::x86_avx512_psrav_q_256
4776 : Intrinsic::x86_avx512_psra_q_256;
4777 else if (
Size ==
'w')
4778 IID = IsImmediate ? Intrinsic::x86_avx2_psrai_w
4779 : Intrinsic::x86_avx2_psra_w;
4784 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_d_512
4785 : IsVariable ? Intrinsic::x86_avx512_psrav_d_512
4786 : Intrinsic::x86_avx512_psra_d_512;
4787 else if (
Size ==
'q')
4788 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_q_512
4789 : IsVariable ? Intrinsic::x86_avx512_psrav_q_512
4790 : Intrinsic::x86_avx512_psra_q_512;
4791 else if (
Size ==
'w')
4792 IID = IsImmediate ? Intrinsic::x86_avx512_psrai_w_512
4793 : Intrinsic::x86_avx512_psra_w_512;
4799 }
else if (Name.starts_with(
"avx512.mask.move.s")) {
4801 }
else if (Name.starts_with(
"avx512.cvtmask2")) {
4803 }
else if (Name.ends_with(
".movntdqa")) {
4807 LoadInst *LI = Builder.CreateAlignedLoad(
4812 }
else if (Name.starts_with(
"fma.vfmadd.") ||
4813 Name.starts_with(
"fma.vfmsub.") ||
4814 Name.starts_with(
"fma.vfnmadd.") ||
4815 Name.starts_with(
"fma.vfnmsub.")) {
4816 bool NegMul = Name[6] ==
'n';
4817 bool NegAcc = NegMul ? Name[8] ==
's' : Name[7] ==
's';
4818 bool IsScalar = NegMul ? Name[12] ==
's' : Name[11] ==
's';
4829 if (NegMul && !IsScalar)
4830 Ops[0] = Builder.CreateFNeg(
Ops[0]);
4831 if (NegMul && IsScalar)
4832 Ops[1] = Builder.CreateFNeg(
Ops[1]);
4834 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4836 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4840 }
else if (Name.starts_with(
"fma4.vfmadd.s")) {
4848 Rep = Builder.CreateIntrinsic(Intrinsic::fma,
Ops[0]->
getType(),
Ops);
4852 }
else if (Name.starts_with(
"avx512.mask.vfmadd.s") ||
4853 Name.starts_with(
"avx512.maskz.vfmadd.s") ||
4854 Name.starts_with(
"avx512.mask3.vfmadd.s") ||
4855 Name.starts_with(
"avx512.mask3.vfmsub.s") ||
4856 Name.starts_with(
"avx512.mask3.vfnmsub.s")) {
4857 bool IsMask3 = Name[11] ==
'3';
4858 bool IsMaskZ = Name[11] ==
'z';
4860 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4861 bool NegMul = Name[2] ==
'n';
4862 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4868 if (NegMul && (IsMask3 || IsMaskZ))
4869 A = Builder.CreateFNeg(
A);
4870 if (NegMul && !(IsMask3 || IsMaskZ))
4871 B = Builder.CreateFNeg(
B);
4873 C = Builder.CreateFNeg(
C);
4875 A = Builder.CreateExtractElement(
A, (
uint64_t)0);
4876 B = Builder.CreateExtractElement(
B, (
uint64_t)0);
4877 C = Builder.CreateExtractElement(
C, (
uint64_t)0);
4884 if (Name.back() ==
'd')
4885 IID = Intrinsic::x86_avx512_vfmadd_f64;
4887 IID = Intrinsic::x86_avx512_vfmadd_f32;
4888 Rep = Builder.CreateIntrinsic(IID,
Ops);
4890 Rep = Builder.CreateFMA(
A,
B,
C);
4899 if (NegAcc && IsMask3)
4904 Rep = Builder.CreateInsertElement(CI->
getArgOperand(IsMask3 ? 2 : 0), Rep,
4906 }
else if (Name.starts_with(
"avx512.mask.vfmadd.p") ||
4907 Name.starts_with(
"avx512.mask.vfnmadd.p") ||
4908 Name.starts_with(
"avx512.mask.vfnmsub.p") ||
4909 Name.starts_with(
"avx512.mask3.vfmadd.p") ||
4910 Name.starts_with(
"avx512.mask3.vfmsub.p") ||
4911 Name.starts_with(
"avx512.mask3.vfnmsub.p") ||
4912 Name.starts_with(
"avx512.maskz.vfmadd.p")) {
4913 bool IsMask3 = Name[11] ==
'3';
4914 bool IsMaskZ = Name[11] ==
'z';
4916 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4917 bool NegMul = Name[2] ==
'n';
4918 bool NegAcc = NegMul ? Name[4] ==
's' : Name[3] ==
's';
4924 if (NegMul && (IsMask3 || IsMaskZ))
4925 A = Builder.CreateFNeg(
A);
4926 if (NegMul && !(IsMask3 || IsMaskZ))
4927 B = Builder.CreateFNeg(
B);
4929 C = Builder.CreateFNeg(
C);
4936 if (Name[Name.size() - 5] ==
's')
4937 IID = Intrinsic::x86_avx512_vfmadd_ps_512;
4939 IID = Intrinsic::x86_avx512_vfmadd_pd_512;
4943 Rep = Builder.CreateFMA(
A,
B,
C);
4951 }
else if (Name.starts_with(
"fma.vfmsubadd.p")) {
4955 if (VecWidth == 128 && EltWidth == 32)
4956 IID = Intrinsic::x86_fma_vfmaddsub_ps;
4957 else if (VecWidth == 256 && EltWidth == 32)
4958 IID = Intrinsic::x86_fma_vfmaddsub_ps_256;
4959 else if (VecWidth == 128 && EltWidth == 64)
4960 IID = Intrinsic::x86_fma_vfmaddsub_pd;
4961 else if (VecWidth == 256 && EltWidth == 64)
4962 IID = Intrinsic::x86_fma_vfmaddsub_pd_256;
4968 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4969 Rep = Builder.CreateIntrinsic(IID,
Ops);
4970 }
else if (Name.starts_with(
"avx512.mask.vfmaddsub.p") ||
4971 Name.starts_with(
"avx512.mask3.vfmaddsub.p") ||
4972 Name.starts_with(
"avx512.maskz.vfmaddsub.p") ||
4973 Name.starts_with(
"avx512.mask3.vfmsubadd.p")) {
4974 bool IsMask3 = Name[11] ==
'3';
4975 bool IsMaskZ = Name[11] ==
'z';
4977 Name = Name.drop_front(IsMask3 || IsMaskZ ? 13 : 12);
4978 bool IsSubAdd = Name[3] ==
's';
4982 if (Name[Name.size() - 5] ==
's')
4983 IID = Intrinsic::x86_avx512_vfmaddsub_ps_512;
4985 IID = Intrinsic::x86_avx512_vfmaddsub_pd_512;
4990 Ops[2] = Builder.CreateFNeg(
Ops[2]);
4992 Rep = Builder.CreateIntrinsic(IID,
Ops);
5001 Value *Odd = Builder.CreateCall(FMA,
Ops);
5002 Ops[2] = Builder.CreateFNeg(
Ops[2]);
5003 Value *Even = Builder.CreateCall(FMA,
Ops);
5009 for (
int i = 0; i != NumElts; ++i)
5010 Idxs[i] = i + (i % 2) * NumElts;
5012 Rep = Builder.CreateShuffleVector(Even, Odd, Idxs);
5020 }
else if (Name.starts_with(
"avx512.mask.pternlog.") ||
5021 Name.starts_with(
"avx512.maskz.pternlog.")) {
5022 bool ZeroMask = Name[11] ==
'z';
5026 if (VecWidth == 128 && EltWidth == 32)
5027 IID = Intrinsic::x86_avx512_pternlog_d_128;
5028 else if (VecWidth == 256 && EltWidth == 32)
5029 IID = Intrinsic::x86_avx512_pternlog_d_256;
5030 else if (VecWidth == 512 && EltWidth == 32)
5031 IID = Intrinsic::x86_avx512_pternlog_d_512;
5032 else if (VecWidth == 128 && EltWidth == 64)
5033 IID = Intrinsic::x86_avx512_pternlog_q_128;
5034 else if (VecWidth == 256 && EltWidth == 64)
5035 IID = Intrinsic::x86_avx512_pternlog_q_256;
5036 else if (VecWidth == 512 && EltWidth == 64)
5037 IID = Intrinsic::x86_avx512_pternlog_q_512;
5043 Rep = Builder.CreateIntrinsic(IID, Args);
5047 }
else if (Name.starts_with(
"avx512.mask.vpmadd52") ||
5048 Name.starts_with(
"avx512.maskz.vpmadd52")) {
5049 bool ZeroMask = Name[11] ==
'z';
5050 bool High = Name[20] ==
'h' || Name[21] ==
'h';
5053 if (VecWidth == 128 && !
High)
5054 IID = Intrinsic::x86_avx512_vpmadd52l_uq_128;
5055 else if (VecWidth == 256 && !
High)
5056 IID = Intrinsic::x86_avx512_vpmadd52l_uq_256;
5057 else if (VecWidth == 512 && !
High)
5058 IID = Intrinsic::x86_avx512_vpmadd52l_uq_512;
5059 else if (VecWidth == 128 &&
High)
5060 IID = Intrinsic::x86_avx512_vpmadd52h_uq_128;
5061 else if (VecWidth == 256 &&
High)
5062 IID = Intrinsic::x86_avx512_vpmadd52h_uq_256;
5063 else if (VecWidth == 512 &&
High)
5064 IID = Intrinsic::x86_avx512_vpmadd52h_uq_512;
5070 Rep = Builder.CreateIntrinsic(IID, Args);
5074 }
else if (Name.starts_with(
"avx512.mask.vpermi2var.") ||
5075 Name.starts_with(
"avx512.mask.vpermt2var.") ||
5076 Name.starts_with(
"avx512.maskz.vpermt2var.")) {
5077 bool ZeroMask = Name[11] ==
'z';
5078 bool IndexForm = Name[17] ==
'i';
5080 }
else if (Name.starts_with(
"avx512.mask.vpdpbusd.") ||
5081 Name.starts_with(
"avx512.maskz.vpdpbusd.") ||
5082 Name.starts_with(
"avx512.mask.vpdpbusds.") ||
5083 Name.starts_with(
"avx512.maskz.vpdpbusds.")) {
5084 bool ZeroMask = Name[11] ==
'z';
5085 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5088 if (VecWidth == 128 && !IsSaturating)
5089 IID = Intrinsic::x86_avx512_vpdpbusd_128;
5090 else if (VecWidth == 256 && !IsSaturating)
5091 IID = Intrinsic::x86_avx512_vpdpbusd_256;
5092 else if (VecWidth == 512 && !IsSaturating)
5093 IID = Intrinsic::x86_avx512_vpdpbusd_512;
5094 else if (VecWidth == 128 && IsSaturating)
5095 IID = Intrinsic::x86_avx512_vpdpbusds_128;
5096 else if (VecWidth == 256 && IsSaturating)
5097 IID = Intrinsic::x86_avx512_vpdpbusds_256;
5098 else if (VecWidth == 512 && IsSaturating)
5099 IID = Intrinsic::x86_avx512_vpdpbusds_512;
5109 if (Args[1]->
getType()->isVectorTy() &&
5112 ->isIntegerTy(32) &&
5113 Args[2]->
getType()->isVectorTy() &&
5116 ->isIntegerTy(32)) {
5117 Type *NewArgType =
nullptr;
5118 if (VecWidth == 128)
5120 else if (VecWidth == 256)
5122 else if (VecWidth == 512)
5128 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5129 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5132 Rep = Builder.CreateIntrinsic(IID, Args);
5136 }
else if (Name.starts_with(
"avx512.mask.vpdpwssd.") ||
5137 Name.starts_with(
"avx512.maskz.vpdpwssd.") ||
5138 Name.starts_with(
"avx512.mask.vpdpwssds.") ||
5139 Name.starts_with(
"avx512.maskz.vpdpwssds.")) {
5140 bool ZeroMask = Name[11] ==
'z';
5141 bool IsSaturating = Name[ZeroMask ? 21 : 20] ==
's';
5144 if (VecWidth == 128 && !IsSaturating)
5145 IID = Intrinsic::x86_avx512_vpdpwssd_128;
5146 else if (VecWidth == 256 && !IsSaturating)
5147 IID = Intrinsic::x86_avx512_vpdpwssd_256;
5148 else if (VecWidth == 512 && !IsSaturating)
5149 IID = Intrinsic::x86_avx512_vpdpwssd_512;
5150 else if (VecWidth == 128 && IsSaturating)
5151 IID = Intrinsic::x86_avx512_vpdpwssds_128;
5152 else if (VecWidth == 256 && IsSaturating)
5153 IID = Intrinsic::x86_avx512_vpdpwssds_256;
5154 else if (VecWidth == 512 && IsSaturating)
5155 IID = Intrinsic::x86_avx512_vpdpwssds_512;
5165 if (Args[1]->
getType()->isVectorTy() &&
5168 ->isIntegerTy(32) &&
5169 Args[2]->
getType()->isVectorTy() &&
5172 ->isIntegerTy(32)) {
5173 Type *NewArgType =
nullptr;
5174 if (VecWidth == 128)
5176 else if (VecWidth == 256)
5178 else if (VecWidth == 512)
5184 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
5185 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
5188 Rep = Builder.CreateIntrinsic(IID, Args);
5192 }
else if (Name ==
"addcarryx.u32" || Name ==
"addcarryx.u64" ||
5193 Name ==
"addcarry.u32" || Name ==
"addcarry.u64" ||
5194 Name ==
"subborrow.u32" || Name ==
"subborrow.u64") {
5196 if (Name[0] ==
'a' && Name.back() ==
'2')
5197 IID = Intrinsic::x86_addcarry_32;
5198 else if (Name[0] ==
'a' && Name.back() ==
'4')
5199 IID = Intrinsic::x86_addcarry_64;
5200 else if (Name[0] ==
's' && Name.back() ==
'2')
5201 IID = Intrinsic::x86_subborrow_32;
5202 else if (Name[0] ==
's' && Name.back() ==
'4')
5203 IID = Intrinsic::x86_subborrow_64;
5210 Value *NewCall = Builder.CreateIntrinsic(IID, Args);
5213 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
5216 Value *CF = Builder.CreateExtractValue(NewCall, 0);
5220 }
else if (Name.starts_with(
"avx512.mask.") &&
5223 }
else if (Name.starts_with(
"bmi.pdep.")) {
5225 }
else if (Name.starts_with(
"bmi.pext.")) {
5235 if (Name.starts_with(
"neon.bfcvt")) {
5236 if (Name.starts_with(
"neon.bfcvtn2")) {
5238 std::iota(LoMask.
begin(), LoMask.
end(), 0);
5240 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5241 Value *Inactive = Builder.CreateShuffleVector(CI->
getOperand(0), LoMask);
5244 return Builder.CreateShuffleVector(Inactive, Trunc, ConcatMask);
5245 }
else if (Name.starts_with(
"neon.bfcvtn")) {
5247 std::iota(ConcatMask.
begin(), ConcatMask.
end(), 0);
5251 dbgs() <<
"Trunc: " << *Trunc <<
"\n";
5252 return Builder.CreateShuffleVector(
5255 return Builder.CreateFPTrunc(CI->
getOperand(0),
5258 }
else if (Name.starts_with(
"sve.fcvt")) {
5261 .
Case(
"sve.fcvt.bf16f32", Intrinsic::aarch64_sve_fcvt_bf16f32_v2)
5262 .
Case(
"sve.fcvtnt.bf16f32",
5263 Intrinsic::aarch64_sve_fcvtnt_bf16f32_v2)
5275 if (Args[1]->
getType() != BadPredTy)
5278 Args[1] = Builder.CreateIntrinsic(Intrinsic::aarch64_sve_convert_to_svbool,
5279 BadPredTy, Args[1]);
5280 Args[1] = Builder.CreateIntrinsic(
5281 Intrinsic::aarch64_sve_convert_from_svbool, GoodPredTy, Args[1]);
5283 return Builder.CreateIntrinsic(NewID, Args,
nullptr,
5287 if (Name ==
"neon.vcvtfp2hf")
5288 return Builder.CreateBitCast(
5289 Builder.CreateFPTrunc(
5293 if (Name ==
"neon.vcvthf2fp")
5294 return Builder.CreateFPExt(
5295 Builder.CreateBitCast(
5305 if (Name ==
"mve.vctp64.old") {
5308 Value *VCTP = Builder.CreateIntrinsic(Intrinsic::arm_mve_vctp64, {},
5311 Value *
C1 = Builder.CreateIntrinsic(
5312 Intrinsic::arm_mve_pred_v2i,
5314 return Builder.CreateIntrinsic(
5315 Intrinsic::arm_mve_pred_i2v,
5317 }
else if (Name ==
"mve.mull.int.predicated.v2i64.v4i32.v4i1" ||
5318 Name ==
"mve.vqdmull.predicated.v2i64.v4i32.v4i1" ||
5319 Name ==
"mve.vldr.gather.base.predicated.v2i64.v2i64.v4i1" ||
5320 Name ==
"mve.vldr.gather.base.wb.predicated.v2i64.v2i64.v4i1" ||
5322 "mve.vldr.gather.offset.predicated.v2i64.p0i64.v2i64.v4i1" ||
5323 Name ==
"mve.vldr.gather.offset.predicated.v2i64.p0.v2i64.v4i1" ||
5324 Name ==
"mve.vstr.scatter.base.predicated.v2i64.v2i64.v4i1" ||
5325 Name ==
"mve.vstr.scatter.base.wb.predicated.v2i64.v2i64.v4i1" ||
5327 "mve.vstr.scatter.offset.predicated.p0i64.v2i64.v2i64.v4i1" ||
5328 Name ==
"mve.vstr.scatter.offset.predicated.p0.v2i64.v2i64.v4i1" ||
5329 Name ==
"cde.vcx1q.predicated.v2i64.v4i1" ||
5330 Name ==
"cde.vcx1qa.predicated.v2i64.v4i1" ||
5331 Name ==
"cde.vcx2q.predicated.v2i64.v4i1" ||
5332 Name ==
"cde.vcx2qa.predicated.v2i64.v4i1" ||
5333 Name ==
"cde.vcx3q.predicated.v2i64.v4i1" ||
5334 Name ==
"cde.vcx3qa.predicated.v2i64.v4i1") {
5335 std::vector<Type *> Tys;
5339 case Intrinsic::arm_mve_mull_int_predicated:
5340 case Intrinsic::arm_mve_vqdmull_predicated:
5341 case Intrinsic::arm_mve_vldr_gather_base_predicated:
5344 case Intrinsic::arm_mve_vldr_gather_base_wb_predicated:
5345 case Intrinsic::arm_mve_vstr_scatter_base_predicated:
5346 case Intrinsic::arm_mve_vstr_scatter_base_wb_predicated:
5350 case Intrinsic::arm_mve_vldr_gather_offset_predicated:
5354 case Intrinsic::arm_mve_vstr_scatter_offset_predicated:
5358 case Intrinsic::arm_cde_vcx1q_predicated:
5359 case Intrinsic::arm_cde_vcx1qa_predicated:
5360 case Intrinsic::arm_cde_vcx2q_predicated:
5361 case Intrinsic::arm_cde_vcx2qa_predicated:
5362 case Intrinsic::arm_cde_vcx3q_predicated:
5363 case Intrinsic::arm_cde_vcx3qa_predicated:
5370 std::vector<Value *>
Ops;
5372 Type *Ty =
Op->getType();
5373 if (Ty->getScalarSizeInBits() == 1) {
5374 Value *
C1 = Builder.CreateIntrinsic(
5375 Intrinsic::arm_mve_pred_v2i,
5377 Op = Builder.CreateIntrinsic(Intrinsic::arm_mve_pred_i2v, {V2I1Ty},
C1);
5382 return Builder.CreateIntrinsic(ID, Tys,
Ops,
nullptr,
5397 auto UpgradeLegacyWMMAIUIntrinsicCall =
5402 Args.push_back(Builder.getFalse());
5406 F->getParent(),
F->getIntrinsicID(), OverloadTys);
5413 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5417 NewCall->copyMetadata(*CI);
5421 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_i32_16x16x64_iu8) {
5422 assert(CI->
arg_size() == 7 &&
"Legacy int_amdgcn_wmma_i32_16x16x64_iu8 "
5423 "intrinsic should have 7 arguments");
5426 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2});
5428 if (
F->getIntrinsicID() == Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8) {
5429 assert(CI->
arg_size() == 8 &&
"Legacy int_amdgcn_swmmac_i32_16x16x128_iu8 "
5430 "intrinsic should have 8 arguments");
5435 return UpgradeLegacyWMMAIUIntrinsicCall(
F, CI, Builder, {
T1, T2, T3, T4});
5438 switch (
F->getIntrinsicID()) {
5441 case Intrinsic::amdgcn_wmma_f32_16x16x4_f32:
5442 case Intrinsic::amdgcn_wmma_f32_16x16x32_bf16:
5443 case Intrinsic::amdgcn_wmma_f32_16x16x32_f16:
5444 case Intrinsic::amdgcn_wmma_f16_16x16x32_f16:
5445 case Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16:
5446 case Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16: {
5461 if (
F->getIntrinsicID() == Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16)
5464 F->getParent(),
F->getIntrinsicID(), Overloads);
5469 auto *NewCall =
cast<CallInst>(Builder.CreateCall(NewDecl, Args, Bundles));
5473 NewCall->copyMetadata(*CI);
5474 NewCall->takeName(CI);
5479 if (Name.starts_with(
"fcmp.") || Name.starts_with(
"icmp.")) {
5485 CallInst *NewCall = Builder.CreateIntrinsicWithoutFolding(
5486 CI->
getType(), Intrinsic::amdgcn_ballot, Cmp);
5494 if (Name.starts_with(
"addrspacecast.nonnull")) {
5497 Value *ASC = Builder.CreateAddrSpaceCast(
5520 if (NumOperands < 3)
5533 bool IsVolatile =
false;
5537 if (NumOperands > 3)
5542 if (NumOperands > 5) {
5544 IsVolatile = !VolatileArg || !VolatileArg->
isZero();
5558 if (VT->getElementType()->isIntegerTy(16)) {
5561 Val = Builder.CreateBitCast(Val, AsBF16);
5569 Builder.CreateAtomicRMW(RMWOp, Ptr, Val, std::nullopt, Order, SSID);
5571 unsigned AddrSpace = PtrTy->getAddressSpace();
5574 RMW->
setMetadata(
"amdgpu.no.fine.grained.memory", EmptyMD);
5576 RMW->
setMetadata(LLVMContext::MD_atomic_ignore_denormal_mode, EmptyMD);
5581 MDNode *RangeNotPrivate =
5584 RMW->
setMetadata(LLVMContext::MD_noalias_addrspace, RangeNotPrivate);
5590 return Builder.CreateBitCast(RMW, RetTy);
5611 return MAV->getMetadata();
5620 if (Name ==
"label") {
5622 }
else if (Name ==
"assign") {
5629 }
else if (Name ==
"declare") {
5633 }
else if (Name ==
"addr") {
5643 unwrapMAVOp(CI, 1), ExprNode,
nullptr,
nullptr,
nullptr);
5644 }
else if (Name ==
"value") {
5647 unsigned ExprOp = 2;
5662 assert(DR &&
"Unhandled intrinsic kind in upgrade to DbgRecord");
5670 int64_t OffsetVal =
Offset->getSExtValue();
5671 return Builder.CreateIntrinsic(OffsetVal >= 0
5672 ? Intrinsic::vector_splice_left
5673 : Intrinsic::vector_splice_right,
5675 {CI->getArgOperand(0), CI->getArgOperand(1),
5676 Builder.getInt32(std::abs(OffsetVal))});
5681 if (Name.starts_with(
"to.fp16")) {
5683 Builder.CreateFPTrunc(CI->
getArgOperand(0), Builder.getHalfTy());
5684 return Builder.CreateBitCast(Cast, CI->
getType());
5687 if (Name.starts_with(
"from.fp16")) {
5689 Builder.CreateBitCast(CI->
getArgOperand(0), Builder.getHalfTy());
5690 return Builder.CreateFPExt(Cast, CI->
getType());
5749 else if (Opcode == Instruction::ICmp)
5752 else if (Opcode == Instruction::FCmp)
5755 else if (Opcode == Instruction::Select)
5760 Rep = Builder.CreateIntrinsic(CI->
getType(), IntrinsicID, Args, {});
5772 if (Defaults.empty())
5775 unsigned OldArgCount = CI->
arg_size();
5776 unsigned NewArgCount = NewFn->
arg_size();
5778 if (OldArgCount < FirstDefault)
5782 if (OldArgCount > NewArgCount)
5787 if (OldArgCount == NewArgCount) {
5799 for (
unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
5800 assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
5801 "missing argument outside the default range");
5802 Type *ParamTy = NewFT->getParamType(Idx);
5807 NewArgs.
push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
5813 CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
5844 if (!Name.consume_front(
"llvm."))
5847 bool IsX86 = Name.consume_front(
"x86.");
5848 bool IsNVVM = Name.consume_front(
"nvvm.");
5849 bool IsAArch64 = Name.consume_front(
"aarch64.");
5850 bool IsARM = Name.consume_front(
"arm.");
5851 bool IsAMDGCN = Name.consume_front(
"amdgcn.");
5852 bool IsDbg = Name.consume_front(
"dbg.");
5854 (Name.consume_front(
"experimental.vector.splice") ||
5855 Name.consume_front(
"vector.splice")) &&
5856 !(Name.starts_with(
".left") || Name.starts_with(
".right"));
5857 Value *Rep =
nullptr;
5859 if (!IsX86 && Name ==
"stackprotectorcheck") {
5861 }
else if (IsNVVM) {
5865 }
else if (IsAArch64) {
5869 }
else if (IsAMDGCN) {
5873 }
else if (IsOldSplice) {
5875 }
else if (Name.consume_front(
"convert.")) {
5877 }
else if (Name ==
"lifetime.start.i64" || Name ==
"lifetime.end.i64") {
5892 const auto &DefaultCase = [&]() ->
void {
5900 "Unknown function for CallBase upgrade and isn't just a name change");
5908 "Return type must have changed");
5909 assert(OldST->getNumElements() ==
5911 "Must have same number of elements");
5914 CallInst *NewCI = Builder.CreateCall(NewFn, Args);
5917 for (
unsigned Idx = 0; Idx < OldST->getNumElements(); ++Idx) {
5918 Value *Elem = Builder.CreateExtractValue(NewCI, Idx);
5919 Res = Builder.CreateInsertValue(Res, Elem, Idx);
5940 case Intrinsic::arm_neon_vst1:
5941 case Intrinsic::arm_neon_vst2:
5942 case Intrinsic::arm_neon_vst3:
5943 case Intrinsic::arm_neon_vst4:
5944 case Intrinsic::arm_neon_vst2lane:
5945 case Intrinsic::arm_neon_vst3lane:
5946 case Intrinsic::arm_neon_vst4lane: {
5948 NewCall = Builder.CreateCall(NewFn, Args);
5951 case Intrinsic::aarch64_sve_bfmlalb_lane_v2:
5952 case Intrinsic::aarch64_sve_bfmlalt_lane_v2:
5953 case Intrinsic::aarch64_sve_bfdot_lane_v2: {
5958 NewCall = Builder.CreateCall(NewFn, Args);
5961 case Intrinsic::aarch64_sve_ld3_sret:
5962 case Intrinsic::aarch64_sve_ld4_sret:
5963 case Intrinsic::aarch64_sve_ld2_sret: {
5971 Name = Name.substr(5);
5978 unsigned MinElts = RetTy->getMinNumElements() /
N;
5980 Value *NewLdCall = Builder.CreateCall(NewFn, Args);
5982 for (
unsigned I = 0;
I <
N;
I++) {
5983 Value *SRet = Builder.CreateExtractValue(NewLdCall,
I);
5984 Ret = Builder.CreateInsertVector(RetTy, Ret, SRet,
I * MinElts);
5990 case Intrinsic::coro_end_async:
5991 case Intrinsic::coro_end: {
5993 if (NewFn->
getIntrinsicID() == Intrinsic::coro_end && Args.size() == 2)
5995 NewCall = Builder.CreateCall(NewFn, Args);
6000 CI->
getModule(), Intrinsic::coro_is_in_ramp);
6001 Value *InRamp = Builder.CreateCall(IsInRamp);
6011 case Intrinsic::vector_extract: {
6013 Name = Name.substr(5);
6014 if (!Name.starts_with(
"aarch64.sve.tuple.get")) {
6019 unsigned MinElts = RetTy->getMinNumElements();
6022 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0), NewIdx});
6026 case Intrinsic::vector_insert: {
6028 Name = Name.substr(5);
6029 if (!Name.starts_with(
"aarch64.sve.tuple")) {
6033 if (Name.starts_with(
"aarch64.sve.tuple.set")) {
6038 NewCall = Builder.CreateCall(
6042 if (Name.starts_with(
"aarch64.sve.tuple.create")) {
6048 assert(
N > 1 &&
"Create is expected to be between 2-4");
6051 unsigned MinElts = RetTy->getMinNumElements() /
N;
6052 for (
unsigned I = 0;
I <
N;
I++) {
6054 Ret = Builder.CreateInsertVector(RetTy, Ret, V,
I * MinElts);
6061 case Intrinsic::arm_neon_bfdot:
6062 case Intrinsic::arm_neon_bfmmla:
6063 case Intrinsic::arm_neon_bfmlalb:
6064 case Intrinsic::arm_neon_bfmlalt:
6065 case Intrinsic::aarch64_neon_bfdot:
6066 case Intrinsic::aarch64_neon_bfmmla:
6067 case Intrinsic::aarch64_neon_bfmlalb:
6068 case Intrinsic::aarch64_neon_bfmlalt: {
6071 "Mismatch between function args and call args");
6072 size_t OperandWidth =
6074 assert((OperandWidth == 64 || OperandWidth == 128) &&
6075 "Unexpected operand width");
6077 auto Iter = CI->
args().begin();
6078 Args.push_back(*Iter++);
6079 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6080 Args.push_back(Builder.CreateBitCast(*Iter++, NewTy));
6081 NewCall = Builder.CreateCall(NewFn, Args);
6085 case Intrinsic::bitreverse:
6086 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6089 case Intrinsic::ctlz:
6090 case Intrinsic::cttz: {
6097 Builder.CreateCall(NewFn, {CI->
getArgOperand(0), Builder.getFalse()});
6101 case Intrinsic::objectsize: {
6102 Value *NullIsUnknownSize =
6106 NewCall = Builder.CreateCall(
6111 case Intrinsic::ctpop:
6112 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(0)});
6114 case Intrinsic::dbg_value: {
6116 Name = Name.substr(5);
6118 if (Name.starts_with(
"dbg.addr")) {
6132 if (
Offset->isNullValue()) {
6133 NewCall = Builder.CreateCall(
6142 case Intrinsic::ptr_annotation:
6150 NewCall = Builder.CreateCall(
6159 case Intrinsic::var_annotation:
6166 NewCall = Builder.CreateCall(
6175 case Intrinsic::riscv_aes32dsi:
6176 case Intrinsic::riscv_aes32dsmi:
6177 case Intrinsic::riscv_aes32esi:
6178 case Intrinsic::riscv_aes32esmi:
6179 case Intrinsic::riscv_sm4ks:
6180 case Intrinsic::riscv_sm4ed: {
6190 Arg0 = Builder.CreateTrunc(Arg0, Builder.getInt32Ty());
6191 Arg1 = Builder.CreateTrunc(Arg1, Builder.getInt32Ty());
6197 NewCall = Builder.CreateCall(NewFn, {Arg0, Arg1, Arg2});
6198 Value *Res = NewCall;
6200 Res = Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6206 case Intrinsic::nvvm_mapa_shared_cluster: {
6210 Value *Res = NewCall;
6211 Res = Builder.CreateAddrSpaceCast(
6218 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster: {
6220 unsigned AS = Args[0]->getType()->getPointerAddressSpace();
6222 Args[0] = Builder.CreateAddrSpaceCast(
6226 Args.push_back(Builder.getInt32(0));
6228 NewCall = Builder.CreateCall(NewFn, Args);
6234 case Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta: {
6239 for (
unsigned I = 0;
I < 4; ++
I)
6241 Args.push_back(Builder.getInt32(0));
6242 Args.push_back(Builder.getInt32(0));
6245 Args.push_back(Builder.getInt1(
false));
6246 Args.push_back(Builder.getInt32(0));
6248 NewCall = Builder.CreateCall(NewFn, Args);
6254 case Intrinsic::nvvm_cp_async_bulk_shared_cta_to_cluster: {
6257 Args[0] = Builder.CreateAddrSpaceCast(
6260 NewCall = Builder.CreateCall(NewFn, Args);
6267#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME) \
6268 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_##ID_SUFFIX:
6270#undef G2S_CLUSTER_CASE
6275 Args[0] = Builder.CreateAddrSpaceCast(
6281 Args.push_back(Builder.getInt32(0));
6283 NewCall = Builder.CreateCall(NewFn, Args);
6290#define G2S_CTA_CASE(ID_SUFFIX, NAME) \
6291 case Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_##ID_SUFFIX:
6299 "expected only the trailing flag_valid_pattern to be missing");
6300 Args.push_back(Builder.getInt32(0));
6302 NewCall = Builder.CreateCall(NewFn, Args);
6308#undef NVVM_TMA_G2S_MODES
6311 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
6312 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
6313 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
6314 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
6315 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
6316 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
6317 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
6318 case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
6320 Name.consume_front(
"llvm.nvvm.cp.async.bulk.tensor.reduce.");
6324 Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
6325 NewCall = Builder.CreateCall(NewFn, Args);
6328 case Intrinsic::nvvm_tcgen05_alloc_cg1:
6329 case Intrinsic::nvvm_tcgen05_alloc_cg2:
6330 case Intrinsic::nvvm_tcgen05_dealloc_cg1:
6331 case Intrinsic::nvvm_tcgen05_dealloc_cg2:
6334 Builder.getFalse()});
6336 case Intrinsic::nvvm_mbarrier_init: {
6340 if (Args.size() == 2)
6341 Args.push_back(Builder.getInt32(0));
6342 NewCall = Builder.CreateCall(NewFn, Args);
6345 case Intrinsic::riscv_sha256sig0:
6346 case Intrinsic::riscv_sha256sig1:
6347 case Intrinsic::riscv_sha256sum0:
6348 case Intrinsic::riscv_sha256sum1:
6349 case Intrinsic::riscv_sm3p0:
6350 case Intrinsic::riscv_sm3p1: {
6357 Builder.CreateTrunc(CI->
getArgOperand(0), Builder.getInt32Ty());
6359 NewCall = Builder.CreateCall(NewFn, Arg);
6361 Builder.CreateIntCast(NewCall, CI->
getType(),
true);
6368 case Intrinsic::x86_xop_vfrcz_ss:
6369 case Intrinsic::x86_xop_vfrcz_sd:
6370 NewCall = Builder.CreateCall(NewFn, {CI->
getArgOperand(1)});
6373 case Intrinsic::x86_xop_vpermil2pd:
6374 case Intrinsic::x86_xop_vpermil2ps:
6375 case Intrinsic::x86_xop_vpermil2pd_256:
6376 case Intrinsic::x86_xop_vpermil2ps_256: {
6380 Args[2] = Builder.CreateBitCast(Args[2], IntIdxTy);
6381 NewCall = Builder.CreateCall(NewFn, Args);
6385 case Intrinsic::x86_sse41_ptestc:
6386 case Intrinsic::x86_sse41_ptestz:
6387 case Intrinsic::x86_sse41_ptestnzc: {
6401 Value *BC0 = Builder.CreateBitCast(Arg0, NewVecTy,
"cast");
6402 Value *BC1 = Builder.CreateBitCast(Arg1, NewVecTy,
"cast");
6404 NewCall = Builder.CreateCall(NewFn, {BC0, BC1});
6408 case Intrinsic::x86_rdtscp: {
6414 NewCall = Builder.CreateCall(NewFn);
6416 Value *
Data = Builder.CreateExtractValue(NewCall, 1);
6419 Value *TSC = Builder.CreateExtractValue(NewCall, 0);
6427 case Intrinsic::x86_sse41_insertps:
6428 case Intrinsic::x86_sse41_dppd:
6429 case Intrinsic::x86_sse41_dpps:
6430 case Intrinsic::x86_sse41_mpsadbw:
6431 case Intrinsic::x86_avx_dp_ps_256:
6432 case Intrinsic::x86_avx2_mpsadbw: {
6438 Args.back() = Builder.CreateTrunc(Args.back(),
Type::getInt8Ty(
C),
"trunc");
6439 NewCall = Builder.CreateCall(NewFn, Args);
6443 case Intrinsic::x86_avx512_mask_cmp_pd_128:
6444 case Intrinsic::x86_avx512_mask_cmp_pd_256:
6445 case Intrinsic::x86_avx512_mask_cmp_pd_512:
6446 case Intrinsic::x86_avx512_mask_cmp_ps_128:
6447 case Intrinsic::x86_avx512_mask_cmp_ps_256:
6448 case Intrinsic::x86_avx512_mask_cmp_ps_512: {
6454 NewCall = Builder.CreateCall(NewFn, Args);
6463 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_128:
6464 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_256:
6465 case Intrinsic::x86_avx512bf16_cvtne2ps2bf16_512:
6466 case Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128:
6467 case Intrinsic::x86_avx512bf16_cvtneps2bf16_256:
6468 case Intrinsic::x86_avx512bf16_cvtneps2bf16_512: {
6472 Intrinsic::x86_avx512bf16_mask_cvtneps2bf16_128)
6473 Args[1] = Builder.CreateBitCast(
6476 NewCall = Builder.CreateCall(NewFn, Args);
6477 Value *Res = Builder.CreateBitCast(
6485 case Intrinsic::x86_avx512bf16_dpbf16ps_128:
6486 case Intrinsic::x86_avx512bf16_dpbf16ps_256:
6487 case Intrinsic::x86_avx512bf16_dpbf16ps_512:{
6491 Args[1] = Builder.CreateBitCast(
6493 Args[2] = Builder.CreateBitCast(
6496 NewCall = Builder.CreateCall(NewFn, Args);
6500 case Intrinsic::thread_pointer: {
6501 NewCall = Builder.CreateCall(NewFn, {});
6505 case Intrinsic::memcpy:
6506 case Intrinsic::memmove:
6507 case Intrinsic::memset: {
6523 NewCall = Builder.CreateCall(NewFn, Args);
6526 C, OldAttrs.getFnAttrs(), OldAttrs.getRetAttrs(),
6527 {OldAttrs.getParamAttrs(0), OldAttrs.getParamAttrs(1),
6528 OldAttrs.getParamAttrs(2), OldAttrs.getParamAttrs(4)});
6533 MemCI->setDestAlignment(
Align->getMaybeAlignValue());
6536 MTI->setSourceAlignment(
Align->getMaybeAlignValue());
6540 case Intrinsic::masked_load:
6541 case Intrinsic::masked_gather:
6542 case Intrinsic::masked_store:
6543 case Intrinsic::masked_scatter: {
6549 auto GetMaybeAlign = [](
Value *
Op) {
6551 uint64_t Val = CI->getZExtValue();
6559 auto GetAlign = [&](
Value *
Op) {
6568 case Intrinsic::masked_load:
6569 NewCall = Builder.CreateMaskedLoad(
6573 case Intrinsic::masked_gather:
6574 NewCall = Builder.CreateMaskedGather(
6580 case Intrinsic::masked_store:
6581 NewCall = Builder.CreateMaskedStore(
6585 case Intrinsic::masked_scatter:
6586 NewCall = Builder.CreateMaskedScatter(
6588 DL.getValueOrABITypeAlignment(
6602 case Intrinsic::lifetime_start:
6603 case Intrinsic::lifetime_end: {
6615 NewCall = Builder.CreateLifetimeStart(Ptr);
6617 NewCall = Builder.CreateLifetimeEnd(Ptr);
6626 case Intrinsic::x86_avx512_vpdpbusd_128:
6627 case Intrinsic::x86_avx512_vpdpbusd_256:
6628 case Intrinsic::x86_avx512_vpdpbusd_512:
6629 case Intrinsic::x86_avx512_vpdpbusds_128:
6630 case Intrinsic::x86_avx512_vpdpbusds_256:
6631 case Intrinsic::x86_avx512_vpdpbusds_512:
6632 case Intrinsic::x86_avx2_vpdpbssd_128:
6633 case Intrinsic::x86_avx2_vpdpbssd_256:
6634 case Intrinsic::x86_avx10_vpdpbssd_512:
6635 case Intrinsic::x86_avx2_vpdpbssds_128:
6636 case Intrinsic::x86_avx2_vpdpbssds_256:
6637 case Intrinsic::x86_avx10_vpdpbssds_512:
6638 case Intrinsic::x86_avx2_vpdpbsud_128:
6639 case Intrinsic::x86_avx2_vpdpbsud_256:
6640 case Intrinsic::x86_avx10_vpdpbsud_512:
6641 case Intrinsic::x86_avx2_vpdpbsuds_128:
6642 case Intrinsic::x86_avx2_vpdpbsuds_256:
6643 case Intrinsic::x86_avx10_vpdpbsuds_512:
6644 case Intrinsic::x86_avx2_vpdpbuud_128:
6645 case Intrinsic::x86_avx2_vpdpbuud_256:
6646 case Intrinsic::x86_avx10_vpdpbuud_512:
6647 case Intrinsic::x86_avx2_vpdpbuuds_128:
6648 case Intrinsic::x86_avx2_vpdpbuuds_256:
6649 case Intrinsic::x86_avx10_vpdpbuuds_512: {
6654 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6655 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6657 NewCall = Builder.CreateCall(NewFn, Args);
6660 case Intrinsic::x86_avx512_vpdpwssd_128:
6661 case Intrinsic::x86_avx512_vpdpwssd_256:
6662 case Intrinsic::x86_avx512_vpdpwssd_512:
6663 case Intrinsic::x86_avx512_vpdpwssds_128:
6664 case Intrinsic::x86_avx512_vpdpwssds_256:
6665 case Intrinsic::x86_avx512_vpdpwssds_512:
6666 case Intrinsic::x86_avx2_vpdpwsud_128:
6667 case Intrinsic::x86_avx2_vpdpwsud_256:
6668 case Intrinsic::x86_avx10_vpdpwsud_512:
6669 case Intrinsic::x86_avx2_vpdpwsuds_128:
6670 case Intrinsic::x86_avx2_vpdpwsuds_256:
6671 case Intrinsic::x86_avx10_vpdpwsuds_512:
6672 case Intrinsic::x86_avx2_vpdpwusd_128:
6673 case Intrinsic::x86_avx2_vpdpwusd_256:
6674 case Intrinsic::x86_avx10_vpdpwusd_512:
6675 case Intrinsic::x86_avx2_vpdpwusds_128:
6676 case Intrinsic::x86_avx2_vpdpwusds_256:
6677 case Intrinsic::x86_avx10_vpdpwusds_512:
6678 case Intrinsic::x86_avx2_vpdpwuud_128:
6679 case Intrinsic::x86_avx2_vpdpwuud_256:
6680 case Intrinsic::x86_avx10_vpdpwuud_512:
6681 case Intrinsic::x86_avx2_vpdpwuuds_128:
6682 case Intrinsic::x86_avx2_vpdpwuuds_256:
6683 case Intrinsic::x86_avx10_vpdpwuuds_512:
6688 Args[1] = Builder.CreateBitCast(Args[1], NewArgType);
6689 Args[2] = Builder.CreateBitCast(Args[2], NewArgType);
6691 NewCall = Builder.CreateCall(NewFn, Args);
6694 assert(NewCall &&
"Should have either set this variable or returned through "
6695 "the default case");
6702 assert(
F &&
"Illegal attempt to upgrade a non-existent intrinsic.");
6716 F->eraseFromParent();
6722 if (NumOperands == 0)
6730 if (NumOperands == 3) {
6734 Metadata *Elts2[] = {ScalarType, ScalarType,
6750 if (NumOperands == 0 || NumOperands % 3 != 0)
6755 for (
unsigned I = 2;
I < NumOperands;
I += 3) {
6760 if (Upgraded ==
Tag)
6770 if (
Opc != Instruction::BitCast)
6774 Type *SrcTy = V->getType();
6791 if (
Opc != Instruction::BitCast)
6794 Type *SrcTy =
C->getType();
6811 if (Flag.getNumOperands() < 3)
6812 return std::nullopt;
6814 return Name->getString();
6815 return std::nullopt;
6821 if (!
getIROptions(M.getContext()).auto_upgrade_debug_info)
6829 if (
NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
6830 auto OpIt =
find_if(ModFlags->operands(), [](
const MDNode *Flag) {
6831 if (auto Name = getModuleFlagNameSafely(*Flag))
6832 return *Name ==
"Debug Info Version";
6835 if (OpIt != ModFlags->op_end()) {
6836 const MDOperand &ValOp = (*OpIt)->getOperand(2);
6843 bool BrokenDebugInfo =
false;
6846 if (!BrokenDebugInfo)
6852 M.getContext().diagnose(Diag);
6859 M.getContext().diagnose(DiagVersion);
6869 StringRef Vect3[3] = {DefaultValue, DefaultValue, DefaultValue};
6872 if (
F->hasFnAttribute(Attr)) {
6875 StringRef S =
F->getFnAttribute(Attr).getValueAsString();
6877 auto [Part, Rest] = S.
split(
',');
6883 const unsigned Dim = DimC -
'x';
6884 assert(Dim < 3 &&
"Unexpected dim char");
6894 F->addFnAttr(Attr, NewAttr);
6898 return S ==
"x" || S ==
"y" || S ==
"z";
6903 if (
K ==
"kernel") {
6915 const unsigned Idx = (AlignIdxValuePair >> 16);
6916 const Align StackAlign =
Align(AlignIdxValuePair & 0xFFFF);
6921 if (
K ==
"maxclusterrank" ||
K ==
"cluster_max_blocks") {
6926 if (
K ==
"minctasm") {
6931 if (
K ==
"maxnreg") {
6936 if (
K.consume_front(
"maxntid") &&
isXYZ(
K)) {
6940 if (
K.consume_front(
"reqntid") &&
isXYZ(
K)) {
6944 if (
K.consume_front(
"cluster_dim_") &&
isXYZ(
K)) {
6948 if (
K ==
"grid_constant") {
6963 NamedMDNode *NamedMD = M.getNamedMetadata(
"nvvm.annotations");
6970 if (!SeenNodes.
insert(MD).second)
6977 assert((MD->getNumOperands() % 2) == 1 &&
"Invalid number of operands");
6984 for (
unsigned j = 1, je = MD->getNumOperands(); j < je; j += 2) {
6986 const MDOperand &V = MD->getOperand(j + 1);
6992 if (NewOperands.
size() > 1)
7005 const char *MarkerKey =
"clang.arc.retainAutoreleasedReturnValueMarker";
7006 NamedMDNode *ModRetainReleaseMarker = M.getNamedMetadata(MarkerKey);
7007 if (ModRetainReleaseMarker) {
7013 ID->getString().split(ValueComp,
"#");
7014 if (ValueComp.
size() == 2) {
7015 std::string NewValue = ValueComp[0].str() +
";" + ValueComp[1].str();
7019 M.eraseNamedMetadata(ModRetainReleaseMarker);
7030 auto UpgradeToIntrinsic = [&](
const char *OldFunc,
7056 bool InvalidCast =
false;
7058 for (
unsigned I = 0, E = CI->
arg_size();
I != E; ++
I) {
7071 Arg = Builder.CreateBitCast(Arg, NewFuncTy->
getParamType(
I));
7073 Args.push_back(Arg);
7080 CallInst *NewCall = Builder.CreateCall(NewFuncTy, NewFn, Args);
7085 Value *NewRetVal = Builder.CreateBitCast(NewCall, CI->
getType());
7098 UpgradeToIntrinsic(
"clang.arc.use", llvm::Intrinsic::objc_clang_arc_use);
7106 std::pair<const char *, llvm::Intrinsic::ID> RuntimeFuncs[] = {
7107 {
"objc_autorelease", llvm::Intrinsic::objc_autorelease},
7108 {
"objc_autoreleasePoolPop", llvm::Intrinsic::objc_autoreleasePoolPop},
7109 {
"objc_autoreleasePoolPush", llvm::Intrinsic::objc_autoreleasePoolPush},
7110 {
"objc_autoreleaseReturnValue",
7111 llvm::Intrinsic::objc_autoreleaseReturnValue},
7112 {
"objc_copyWeak", llvm::Intrinsic::objc_copyWeak},
7113 {
"objc_destroyWeak", llvm::Intrinsic::objc_destroyWeak},
7114 {
"objc_initWeak", llvm::Intrinsic::objc_initWeak},
7115 {
"objc_loadWeak", llvm::Intrinsic::objc_loadWeak},
7116 {
"objc_loadWeakRetained", llvm::Intrinsic::objc_loadWeakRetained},
7117 {
"objc_moveWeak", llvm::Intrinsic::objc_moveWeak},
7118 {
"objc_release", llvm::Intrinsic::objc_release},
7119 {
"objc_retain", llvm::Intrinsic::objc_retain},
7120 {
"objc_retainAutorelease", llvm::Intrinsic::objc_retainAutorelease},
7121 {
"objc_retainAutoreleaseReturnValue",
7122 llvm::Intrinsic::objc_retainAutoreleaseReturnValue},
7123 {
"objc_retainAutoreleasedReturnValue",
7124 llvm::Intrinsic::objc_retainAutoreleasedReturnValue},
7125 {
"objc_retainBlock", llvm::Intrinsic::objc_retainBlock},
7126 {
"objc_storeStrong", llvm::Intrinsic::objc_storeStrong},
7127 {
"objc_storeWeak", llvm::Intrinsic::objc_storeWeak},
7128 {
"objc_unsafeClaimAutoreleasedReturnValue",
7129 llvm::Intrinsic::objc_unsafeClaimAutoreleasedReturnValue},
7130 {
"objc_retainedObject", llvm::Intrinsic::objc_retainedObject},
7131 {
"objc_unretainedObject", llvm::Intrinsic::objc_unretainedObject},
7132 {
"objc_unretainedPointer", llvm::Intrinsic::objc_unretainedPointer},
7133 {
"objc_retain_autorelease", llvm::Intrinsic::objc_retain_autorelease},
7134 {
"objc_sync_enter", llvm::Intrinsic::objc_sync_enter},
7135 {
"objc_sync_exit", llvm::Intrinsic::objc_sync_exit},
7136 {
"objc_arc_annotation_topdown_bbstart",
7137 llvm::Intrinsic::objc_arc_annotation_topdown_bbstart},
7138 {
"objc_arc_annotation_topdown_bbend",
7139 llvm::Intrinsic::objc_arc_annotation_topdown_bbend},
7140 {
"objc_arc_annotation_bottomup_bbstart",
7141 llvm::Intrinsic::objc_arc_annotation_bottomup_bbstart},
7142 {
"objc_arc_annotation_bottomup_bbend",
7143 llvm::Intrinsic::objc_arc_annotation_bottomup_bbend}};
7145 for (
auto &
I : RuntimeFuncs)
7146 UpgradeToIntrinsic(
I.first,
I.second);
7170 std::optional<bool> UseAddressDisc;
7173 if (
const NamedMDNode *ModFlags = M.getModuleFlagsMetadata()) {
7174 for (
const MDNode *Flag : ModFlags->operands()) {
7176 if (Name && (*Name ==
"ptrauth-init-fini" ||
7177 *Name ==
"ptrauth-init-fini-address-discrimination"))
7182 auto UpgradeSinglePointer = [&UseAddressDisc](
Constant *CV) ->
Constant * {
7183 constexpr unsigned ExpectedConstDisc = 0xD9D4;
7184 constexpr unsigned ExpectedAddressMarker = 1;
7187 if (!CPA || !CPA->getDiscriminator()->equalsInt(ExpectedConstDisc))
7190 bool HasAddressDisc;
7191 if (!CPA->hasAddressDiscriminator())
7192 HasAddressDisc =
false;
7193 else if (CPA->hasSpecialAddressDiscriminator(ExpectedAddressMarker))
7194 HasAddressDisc =
true;
7198 if (UseAddressDisc && *UseAddressDisc != HasAddressDisc)
7201 UseAddressDisc = HasAddressDisc;
7202 return CPA->getPointer();
7206 using PendingUpgrade = std::pair<GlobalVariable *, Constant *>;
7209 for (
const char *Name : {
"llvm.global_ctors",
"llvm.global_dtors"}) {
7211 if (!GV || !GV->hasInitializer())
7215 if (!OldStructorsArray || OldStructorsArray->getNumOperands() == 0)
7218 std::vector<Constant *> NewStructors;
7219 NewStructors.reserve(OldStructorsArray->getNumOperands());
7221 for (
Use &U : OldStructorsArray->operands()) {
7230 Func = UpgradeSinglePointer(Func);
7234 NewStructors.push_back(
7243 if (GlobalArraysToUpgrade.
empty())
7245 assert(UseAddressDisc.has_value());
7247 for (
auto [GV, NewInit] : GlobalArraysToUpgrade)
7248 GV->setInitializer(NewInit);
7251 M.addModuleFlag(
Module::Error,
"ptrauth-init-fini-address-discrimination",
7261 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7265 bool HasObjCFlag =
false, HasClassProperties =
false;
7266 bool HasSwiftVersionFlag =
false;
7267 uint8_t SwiftMajorVersion, SwiftMinorVersion;
7274 if (
Op->getNumOperands() != 3)
7288 if (ID->getString() ==
"Objective-C Image Info Version")
7290 if (ID->getString() ==
"Objective-C Class Properties")
7291 HasClassProperties =
true;
7293 if (ID->getString() ==
"PIC Level") {
7294 if (
auto *Behavior =
7296 uint64_t V = Behavior->getLimitedValue();
7302 if (ID->getString() ==
"PIE Level")
7303 if (
auto *Behavior =
7311 if (ID->getString() ==
"branch-target-enforcement" ||
7312 (ID->getString().starts_with(
"sign-return-address") &&
7313 ID->getString() !=
"sign-return-address-harden")) {
7314 if (
auto *Behavior =
7320 Op->getOperand(1),
Op->getOperand(2)};
7330 if (ID->getString() ==
"Objective-C Image Info Section") {
7333 Value->getString().split(ValueComp,
" ");
7334 if (ValueComp.
size() != 1) {
7335 std::string NewValue;
7336 for (
auto &S : ValueComp)
7337 NewValue += S.str();
7348 if (ID->getString() ==
"Objective-C Garbage Collection") {
7351 assert(Md->getValue() &&
"Expected non-empty metadata");
7352 auto Type = Md->getValue()->getType();
7355 unsigned Val = Md->getValue()->getUniqueInteger().getZExtValue();
7356 if ((Val & 0xff) != Val) {
7357 HasSwiftVersionFlag =
true;
7358 SwiftABIVersion = (Val & 0xff00) >> 8;
7359 SwiftMajorVersion = (Val & 0xff000000) >> 24;
7360 SwiftMinorVersion = (Val & 0xff0000) >> 16;
7371 if (ID->getString() ==
"amdgpu_code_object_version") {
7374 MDString::get(M.getContext(),
"amdhsa_code_object_version"),
7383 if (M.getTargetTriple().isPPC() && ID->getString() ==
"float-abi") {
7412 if (HasObjCFlag && !HasClassProperties) {
7418 if (HasSwiftVersionFlag) {
7422 ConstantInt::get(Int8Ty, SwiftMajorVersion));
7424 ConstantInt::get(Int8Ty, SwiftMinorVersion));
7432 NamedMDNode *CFIConsts = M.getNamedMetadata(
"cfi.functions");
7436 auto MatchesVersion = [](
const MDNode *
Op) {
7437 return Op->getNumOperands() >= 3 &&
7451 assert(!MatchesVersion(
Op) &&
"Unexpected mix of CFIConstant formats");
7452 assert(
Op->getNumOperands() >= 2 &&
7453 "Expected at least 2 operands - name and linkage type");
7465 for (
unsigned J = 2, EJ =
Op->getNumOperands(); J != EJ; ++J)
7476 auto TrimSpaces = [](
StringRef Section) -> std::string {
7478 Section.split(Components,
',');
7483 for (
auto Component : Components)
7484 OS <<
',' << Component.trim();
7489 for (
auto &GV : M.globals()) {
7490 if (!GV.hasSection())
7495 if (!Section.starts_with(
"__DATA, __objc_catlist"))
7500 GV.setSection(TrimSpaces(Section));
7516struct StrictFPUpgradeVisitor :
public InstVisitor<StrictFPUpgradeVisitor> {
7517 StrictFPUpgradeVisitor() =
default;
7520 if (!
Call.isStrictFP())
7526 Call.removeFnAttr(Attribute::StrictFP);
7527 Call.addFnAttr(Attribute::NoBuiltin);
7532struct AMDGPUUnsafeFPAtomicsUpgradeVisitor
7533 :
public InstVisitor<AMDGPUUnsafeFPAtomicsUpgradeVisitor> {
7534 AMDGPUUnsafeFPAtomicsUpgradeVisitor() =
default;
7536 void visitAtomicRMWInst(AtomicRMWInst &RMW) {
7551 if (!
F.isDeclaration() && !
F.hasFnAttribute(Attribute::StrictFP)) {
7552 StrictFPUpgradeVisitor SFPV;
7557 F.removeRetAttrs(AttributeFuncs::typeIncompatible(
7558 F.getReturnType(),
F.getAttributes().getRetAttrs()));
7559 for (
auto &Arg :
F.args())
7561 AttributeFuncs::typeIncompatible(Arg.getType(), Arg.getAttributes()));
7563 bool AddingAttrs =
false, RemovingAttrs =
false;
7564 AttrBuilder AttrsToAdd(
F.getContext());
7569 if (
Attribute A =
F.getFnAttribute(
"implicit-section-name");
7570 A.isValid() &&
A.isStringAttribute()) {
7571 F.setSection(
A.getValueAsString());
7573 RemovingAttrs =
true;
7577 A.isValid() &&
A.isStringAttribute()) {
7580 AddingAttrs = RemovingAttrs =
true;
7583 if (
Attribute A =
F.getFnAttribute(
"uniform-work-group-size");
7584 A.isValid() &&
A.isStringAttribute() && !
A.getValueAsString().empty()) {
7586 RemovingAttrs =
true;
7587 if (
A.getValueAsString() ==
"true") {
7588 AttrsToAdd.addAttribute(
"uniform-work-group-size");
7597 if (
Attribute A =
F.getFnAttribute(
"amdgpu-unsafe-fp-atomics");
7600 if (
A.getValueAsBool()) {
7601 AMDGPUUnsafeFPAtomicsUpgradeVisitor Visitor;
7607 AttrsToRemove.
addAttribute(
"amdgpu-unsafe-fp-atomics");
7608 RemovingAttrs =
true;
7615 bool HandleDenormalMode =
false;
7617 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math"); Attr.isValid()) {
7620 DenormalFPMath = ParsedMode;
7622 AddingAttrs = RemovingAttrs =
true;
7623 HandleDenormalMode =
true;
7627 if (
Attribute Attr =
F.getFnAttribute(
"denormal-fp-math-f32");
7631 DenormalFPMathF32 = ParsedMode;
7633 AddingAttrs = RemovingAttrs =
true;
7634 HandleDenormalMode =
true;
7638 if (HandleDenormalMode)
7639 AttrsToAdd.addDenormalFPEnvAttr(
7643 F.removeFnAttrs(AttrsToRemove);
7646 F.addFnAttrs(AttrsToAdd);
7652 if (!
F.hasFnAttribute(FnAttrName)) {
7653 F.addFnAttr(FnAttrName,
Value);
7655 <<
"\", function: " <<
F.getName() <<
"\n");
7663 if (!
F.hasFnAttribute(FnAttrName)) {
7665 F.addFnAttr(FnAttrName);
7667 <<
", function: " <<
F.getName() <<
"\n");
7670 auto A =
F.getFnAttribute(FnAttrName);
7671 if (
"false" ==
A.getValueAsString()) {
7672 F.removeFnAttr(FnAttrName);
7674 <<
"=\"false\", function: " <<
F.getName() <<
"\n");
7675 }
else if (
"true" ==
A.getValueAsString()) {
7676 F.removeFnAttr(FnAttrName);
7677 F.addFnAttr(FnAttrName);
7679 <<
"=\"true\", function: " <<
F.getName() <<
"\n");
7686 M.setModuleFlag(Behavior,
Key, Val);
7687 LLVM_DEBUG(
dbgs() <<
"Converted module flag: " <<
"{" << Behavior <<
", "
7688 <<
Key <<
", " << Val <<
"}\n");
7692 Triple T(M.getTargetTriple());
7693 if (!
T.isThumb() && !
T.isARM() && !
T.isAArch64())
7696 uint64_t BTEValue = 0;
7697 uint64_t BPPLRValue = 0;
7698 uint64_t GCSValue = 0;
7699 uint64_t SRAValue = 0;
7700 uint64_t SRAALLValue = 0;
7701 uint64_t SRABKeyValue = 0;
7703 NamedMDNode *ModFlags = M.getModuleFlagsMetadata();
7707 if (
Op->getNumOperands() != 3)
7716 uint64_t *ValPtr = IDStr ==
"branch-target-enforcement" ? &BTEValue
7717 : IDStr ==
"branch-protection-pauth-lr" ? &BPPLRValue
7718 : IDStr ==
"guarded-control-stack" ? &GCSValue
7719 : IDStr ==
"sign-return-address" ? &SRAValue
7720 : IDStr ==
"sign-return-address-all" ? &SRAALLValue
7721 : IDStr ==
"sign-return-address-with-bkey"
7727 *ValPtr = CI->getZExtValue();
7731 LLVM_DEBUG(
dbgs() <<
"Found module flag: " << IDStr <<
"(" << *ValPtr
7736 bool BTE = BTEValue == 1;
7737 bool BPPLR = BPPLRValue == 1;
7738 bool GCS = GCSValue == 1;
7739 bool SRA = SRAValue == 1;
7742 if (SRA && SRAALLValue == 1)
7743 SignTypeValue =
"all";
7746 if (SRA && SRABKeyValue == 1)
7747 SignKeyValue =
"b_key";
7749 for (
Function &
F : M.getFunctionList()) {
7750 if (
F.isDeclaration())
7757 if (
auto A =
F.getFnAttribute(
"sign-return-address");
7758 A.isValid() &&
"none" ==
A.getValueAsString()) {
7759 F.removeFnAttr(
"sign-return-address");
7760 F.removeFnAttr(
"sign-return-address-key");
7776 if (SRAALLValue == 1)
7778 if (SRABKeyValue == 1) {
7807 if (
T->getNumOperands() < 1)
7812 if (S->getString().starts_with(
"llvm.vectorizer."))
7818 StringRef OldPrefix =
"llvm.vectorizer.";
7821 if (OldTag ==
"llvm.vectorizer.unroll")
7833 if (
T->getNumOperands() < 1)
7845 if (!OldTag->getString().starts_with(
"llvm.vectorizer."))
7858 Ops.reserve(
T->getNumOperands());
7859 Ops.push_back(NewTag);
7860 for (
unsigned I = 1,
E =
T->getNumOperands();
I !=
E; ++
I)
7861 Ops.push_back(
T->getOperand(
I));
7878 if (
T->isDistinct()) {
7879 for (
unsigned I = 0, E =
T->getNumOperands();
I < E; ++
I) {
7891 Ops.reserve(
T->getNumOperands());
7902 if ((
T.isSPIR() || (
T.isSPIRV() && !
T.isSPIRVLogical())) &&
7903 !
DL.contains(
"-G") && !
DL.starts_with(
"G")) {
7904 return DL.empty() ? std::string(
"G1") : (
DL +
"-G1").str();
7907 if (
T.isLoongArch64() ||
T.isRISCV64()) {
7909 auto I =
DL.find(
"-n64-");
7911 return (
DL.take_front(
I) +
"-n32:64-" +
DL.drop_front(
I + 5)).str();
7916 std::string Res =
DL.str();
7919 if (!
DL.contains(
"-G") && !
DL.starts_with(
"G"))
7920 Res.append(Res.empty() ?
"G1" :
"-G1");
7928 if (!
DL.contains(
"-ni") && !
DL.starts_with(
"ni"))
7929 Res.append(
"-ni:7:8:9");
7931 if (
DL.ends_with(
"ni:7"))
7933 if (
DL.ends_with(
"ni:7:8"))
7938 if (!
DL.contains(
"-p7") && !
DL.starts_with(
"p7"))
7939 Res.append(
"-p7:160:256:256:32");
7940 if (!
DL.contains(
"-p8") && !
DL.starts_with(
"p8"))
7941 Res.append(
"-p8:128:128:128:48");
7942 constexpr StringRef OldP8(
"-p8:128:128-");
7943 if (
DL.contains(OldP8))
7944 Res.replace(Res.find(OldP8), OldP8.
size(),
"-p8:128:128:128:48-");
7945 if (!
DL.contains(
"-p9") && !
DL.starts_with(
"p9"))
7946 Res.append(
"-p9:192:256:256:32");
7951 for (
StringRef AS : {
"p10",
"p11",
"p12",
"p13",
"p14",
"p15"}) {
7952 if (!
DL.contains((
"-" + AS).str()) && !
DL.starts_with(AS))
7953 Res.append((
"-" + AS +
":32:32").str());
7958 if (!
DL.contains(
"m:e"))
7959 Res = Res.empty() ?
"m:e" :
"m:e-" + Res;
7964 if (
T.isSystemZ() && !
DL.empty()) {
7966 if (!
DL.contains(
"-S64"))
7967 return "E-S64" +
DL.drop_front(1).str();
7971 auto AddPtr32Ptr64AddrSpaces = [&
DL, &Res]() {
7974 StringRef AddrSpaces{
"-p270:32:32-p271:32:32-p272:64:64"};
7975 if (!
DL.contains(AddrSpaces)) {
7977 Regex R(
"^([Ee]-m:[a-z](-p:32:32)?)(-.*)$");
7978 if (R.match(Res, &
Groups))
7984 if (
T.isAArch64()) {
7986 if (!
DL.empty() && !
DL.contains(
"-Fn32"))
7987 Res.append(
"-Fn32");
7988 AddPtr32Ptr64AddrSpaces();
7992 if (
T.isSPARC() || (
T.isMIPS64() && !
DL.contains(
"m:m")) ||
T.isPPC64() ||
7996 std::string I64 =
"-i64:64";
7997 std::string I128 =
"-i128:128";
7999 size_t Pos = Res.find(I64);
8000 if (Pos !=
size_t(-1))
8001 Res.insert(Pos + I64.size(), I128);
8005 if (
T.isPPC() &&
T.isOSAIX() && !
DL.contains(
"f64:32:64") && !
DL.empty()) {
8006 size_t Pos = Res.find(
"-S128");
8009 Res.insert(Pos,
"-f64:32:64");
8014 if (
T.isARM() && !
DL.empty() && !
DL.contains(
"Fi") && !
DL.contains(
"Fn")) {
8016 size_t Pos = Res.
find(p3232);
8018 Res.insert(Pos + p3232.
size(),
"-Fi8");
8024 AddPtr32Ptr64AddrSpaces();
8032 if (!
T.isOSIAMCU()) {
8033 std::string I128 =
"-i128:128";
8036 Regex R(
"^(e(-[mpi][^-]*)*)((-[^mpi][^-]*)*)$");
8037 if (R.match(Res, &
Groups))
8045 if (
T.isWindowsMSVCEnvironment() && !
T.isArch64Bit()) {
8047 auto I =
Ref.find(
"-f80:32-");
8049 Res = (
Ref.take_front(
I) +
"-f80:128-" +
Ref.drop_front(
I + 8)).str();
8057 Attribute A =
B.getAttribute(
"no-frame-pointer-elim");
8060 FramePointer =
A.getValueAsString() ==
"true" ?
"all" :
"none";
8061 B.removeAttribute(
"no-frame-pointer-elim");
8063 if (
B.contains(
"no-frame-pointer-elim-non-leaf")) {
8065 if (FramePointer !=
"all")
8066 FramePointer =
"non-leaf";
8067 B.removeAttribute(
"no-frame-pointer-elim-non-leaf");
8069 if (!FramePointer.
empty())
8070 B.addAttribute(
"frame-pointer", FramePointer);
8072 A =
B.getAttribute(
"null-pointer-is-valid");
8075 bool NullPointerIsValid =
A.getValueAsString() ==
"true";
8076 B.removeAttribute(
"null-pointer-is-valid");
8077 if (NullPointerIsValid)
8078 B.addAttribute(Attribute::NullPointerIsValid);
8081 A =
B.getAttribute(
"uniform-work-group-size");
8085 bool IsTrue = Val ==
"true";
8086 B.removeAttribute(
"uniform-work-group-size");
8088 B.addAttribute(
"uniform-work-group-size");
8099 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 Intrinsic::ID shouldUpgradeNVPTXTcgen05AllocDeallocIntrinsic(Function *F, StringRef Name)
static std::optional< StringRef > getModuleFlagNameSafely(const MDNode &Flag)
static bool consumeNVVMPtrAddrSpace(StringRef &Name)
static Metadata * makeBooleanLoopNode(LLVMContext &C, const BooleanLoopTags &Tags, const MDOperand &Op)
Build the single-operand node that replaces a boolean operand: nonzero selects the enable tag,...
#define G2S_CLUSTER_CASE(ID_SUFFIX, NAME)
static bool shouldUpgradeX86Intrinsic(Function *F, StringRef Name)
static Value * upgradeX86PSRLDQIntrinsics(IRBuilder<> &Builder, Value *Op, unsigned Shift)
static unsigned getFunctionalOpcodeForVP(StringRef Name)
static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F, StringRef Name, SmallVectorImpl< Type * > &OvlTys)
static Intrinsic::ID shouldUpgradeNVPTXTcgen05CommitSharedIntrinsic(Function *F, StringRef Name)
static 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 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 bool isOverloaded(ID id)
Returns true if the intrinsic can be overloaded.
LLVM_ABI FunctionType * getType(LLVMContext &Context, ID id, ArrayRef< Type * > OverloadTys={})
Return the function type for an intrinsic.
LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT, SmallVectorImpl< Type * > &OverloadTys, raw_ostream &OS=nulls())
Returns true if FT is a valid function type for intrinsic ID.
LLVM_ABI bool hasStructReturnType(ID id)
Returns true if id has a struct return type.
LLVM_ABI std::pair< unsigned, ArrayRef< uint64_t > > getAllDefaultArgValues(ID IID)
Returns the first default argument index and an ArrayRef of all default values for the trailing param...
@ ADDRESS_SPACE_SHARED_CLUSTER
constexpr StringLiteral GridConstant("nvvm.grid_constant")
constexpr StringLiteral MaxNTID("nvvm.maxntid")
constexpr StringLiteral MaxNReg("nvvm.maxnreg")
constexpr StringLiteral MinCTASm("nvvm.minctasm")
constexpr StringLiteral ReqNTID("nvvm.reqntid")
constexpr StringLiteral MaxClusterRank("nvvm.maxclusterrank")
constexpr StringLiteral ClusterDim("nvvm.cluster_dim")
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract_or_null(Y &&MD)
Extract a Value from Metadata, if any, allowing null.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, bool > hasa(Y &&MD)
Check whether Metadata has a Value.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > dyn_extract(Y &&MD)
Extract a Value from Metadata, if any.
std::enable_if_t< detail::IsValidPointer< X, Y >::value, X * > extract(Y &&MD)
Extract a Value from Metadata.
This is an optimization pass for GlobalISel generic memory operations.
LLVM_ABI void UpgradeIntrinsicCall(CallBase *CB, Function *NewFn)
This is the complement to the above, replacing a specific call to an intrinsic function with a call t...
LLVM_ABI void UpgradeSectionAttributes(Module &M)
auto size(R &&Range, std::enable_if_t< std::is_base_of< std::random_access_iterator_tag, typename std::iterator_traits< decltype(Range.begin())>::iterator_category >::value, void > *=nullptr)
Get the size of a range.
LLVM_ABI void UpgradeInlineAsmString(std::string *AsmStr)
Upgrade comment in call to inline asm that represents an objc retain release marker.
bool isValidAtomicOrdering(Int I)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
@ Load
The value being inserted comes from a load (InsertElement only).
StringRef getLongDoubleFormatName(LongDoubleFormat Format)
Returns the IR floating-point type name for a LongDoubleFormat.
LongDoubleFormat
The floating-point format used for the target's "long double" type.
LLVM_ABI bool UpgradeIntrinsicFunction(Function *F, Function *&NewFn, bool CanUpgradeDebugIntrinsicsToRecords=true)
This is a more granular function that simply checks an intrinsic function for upgrading,...
LLVM_ABI MDNode * upgradeInstructionLoopAttachment(MDNode &N)
Upgrade the loop attachment metadata node.
auto dyn_cast_if_present(const Y &Val)
dyn_cast_if_present<X> - Functionally identical to dyn_cast, except that a null (or none in the case ...
LLVM_ABI void UpgradeAttributes(AttrBuilder &B)
Upgrade attributes that changed format or kind.
LLVM_ABI void UpgradeCallsToIntrinsic(Function *F)
This is an auto-upgrade hook for any old intrinsic function syntaxes which need to have both the func...
LLVM_ABI void UpgradeNVVMAnnotations(Module &M)
Convert legacy nvvm.annotations metadata to appropriate function attributes.
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
LLVM_ABI bool UpgradeModuleFlags(Module &M)
This checks for module flags which should be upgraded.
std::string utostr(uint64_t X, bool isNeg=false)
constexpr bool isPowerOf2_64(uint64_t Value)
Return true if the argument is a power of two > 0 (64 bit edition.)
LLVM_ABI bool UpgradeCFIFunctionsMetadata(Module &M)
Upgrade the cfi.functions metadata node by calculating and inserting the GUID for each function entry...
LLVM_ABI void copyModuleAttrToFunctions(Module &M)
Copies module attributes to the functions in the module.
LLVM_ABI void UpgradeOperandBundles(std::vector< OperandBundleDef > &OperandBundles)
Upgrade operand bundles (without knowing about their user instruction).
LLVM_ABI Constant * UpgradeBitCastExpr(unsigned Opc, Constant *C, Type *DestTy)
This is an auto-upgrade for bitcast constant expression between pointers with different address space...
auto dyn_cast_or_null(const Y &Val)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI std::string UpgradeDataLayoutString(StringRef DL, StringRef Triple)
Upgrade the datalayout string by adding a section for address space pointers.
bool none_of(R &&Range, UnaryPredicate P)
Provide wrappers to std::none_of which take ranges instead of having to pass begin/end explicitly.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
LLVM_ABI MDNode * UpgradeTBAAStructNode(MDNode &TBAAStructNode)
If the given !tbaa.struct node has old-style scalar field tags, return an equivalent node with each f...
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
LLVM_ABI GlobalVariable * UpgradeGlobalVariable(GlobalVariable *GV)
This checks for global variables which should be upgraded.
LLVM_ATTRIBUTE_VISIBILITY_DEFAULT AnalysisKey InnerAnalysisManagerProxy< AnalysisManagerT, IRUnitT, ExtraArgTs... >::Key
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
LLVM_ABI bool StripDebugInfo(Module &M)
Strip debug info in the module if it exists.
auto drop_end(T &&RangeOrContainer, size_t N=1)
Return a range covering RangeOrContainer with the last N elements excluded.
AtomicOrdering
Atomic ordering for LLVM's memory model.
@ Ref
The access may reference the value stored in memory.
std::string join(IteratorT Begin, IteratorT End, StringRef Separator)
Joins the strings in the range [Begin, End), adding Separator between the elements.
const BooleanLoopTags * findBooleanLoopTags(StringRef Name)
Return the replacement tags for the enable tag Name, or nullptr.
OperandBundleDefT< Value * > OperandBundleDef
LLVM_ABI Instruction * UpgradeBitCastInst(unsigned Opc, Value *V, Type *DestTy, Instruction *&Temp)
This is an auto-upgrade for bitcast between pointers with different address spaces: the instruction i...
DWARFExpression::Operation Op
RoundingMode
Rounding mode.
@ TowardZero
roundTowardZero.
@ NearestTiesToEven
roundTiesToEven.
@ Dynamic
Denotes mode unknown at compile time.
@ TowardPositive
roundTowardPositive.
@ TowardNegative
roundTowardNegative.
ArrayRef(const T &OneElt) -> ArrayRef< T >
const IROptions & getIROptions(const LLVMContext &C)
DenormalMode parseDenormalFPAttribute(StringRef Str)
Returns the denormal mode to use for inputs and outputs.
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
auto find_if(R &&Range, UnaryPredicate P)
Provide wrappers to std::find_if which take ranges instead of having to pass begin/end explicitly.
void erase_if(Container &C, UnaryPredicate P)
Provide a container algorithm similar to C++ Library Fundamentals v2's erase_if which is equivalent t...
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.