From 6c9156717e226a64639b79428febd519a1cbb5d6 Mon Sep 17 00:00:00 2001 From: Guillaume Gomez Date: Wed, 13 Mar 2024 16:44:54 +0100 Subject: [PATCH] Regen intrinsics conversions --- src/intrinsic/archs.rs | 33 +++++++++++++++++++++++++++++++++ 1 file changed, 33 insertions(+) diff --git a/src/intrinsic/archs.rs b/src/intrinsic/archs.rs index c4ae1751fa0..f7500933789 100644 --- a/src/intrinsic/archs.rs +++ b/src/intrinsic/archs.rs @@ -74,6 +74,10 @@ match name { "llvm.amdgcn.cvt.sr.bf8.f32" => "__builtin_amdgcn_cvt_sr_bf8_f32", "llvm.amdgcn.cvt.sr.fp8.f32" => "__builtin_amdgcn_cvt_sr_fp8_f32", "llvm.amdgcn.dispatch.id" => "__builtin_amdgcn_dispatch_id", + "llvm.amdgcn.dot4.f32.bf8.bf8" => "__builtin_amdgcn_dot4_f32_bf8_bf8", + "llvm.amdgcn.dot4.f32.bf8.fp8" => "__builtin_amdgcn_dot4_f32_bf8_fp8", + "llvm.amdgcn.dot4.f32.fp8.bf8" => "__builtin_amdgcn_dot4_f32_fp8_bf8", + "llvm.amdgcn.dot4.f32.fp8.fp8" => "__builtin_amdgcn_dot4_f32_fp8_fp8", "llvm.amdgcn.ds.add.gs.reg.rtn" => "__builtin_amdgcn_ds_add_gs_reg_rtn", "llvm.amdgcn.ds.bpermute" => "__builtin_amdgcn_ds_bpermute", "llvm.amdgcn.ds.fadd.v2bf16" => "__builtin_amdgcn_ds_atomic_fadd_v2bf16", @@ -2291,6 +2295,10 @@ match name { "llvm.loongarch.csrxchg.d" => "__builtin_loongarch_csrxchg_d", "llvm.loongarch.csrxchg.w" => "__builtin_loongarch_csrxchg_w", "llvm.loongarch.dbar" => "__builtin_loongarch_dbar", + "llvm.loongarch.frecipe.d" => "__builtin_loongarch_frecipe_d", + "llvm.loongarch.frecipe.s" => "__builtin_loongarch_frecipe_s", + "llvm.loongarch.frsqrte.d" => "__builtin_loongarch_frsqrte_d", + "llvm.loongarch.frsqrte.s" => "__builtin_loongarch_frsqrte_s", "llvm.loongarch.ibar" => "__builtin_loongarch_ibar", "llvm.loongarch.iocsrrd.b" => "__builtin_loongarch_iocsrrd_b", "llvm.loongarch.iocsrrd.d" => "__builtin_loongarch_iocsrrd_d", @@ -2529,6 +2537,8 @@ match name { "llvm.loongarch.lasx.xvfnmsub.s" => "__builtin_lasx_xvfnmsub_s", "llvm.loongarch.lasx.xvfrecip.d" => "__builtin_lasx_xvfrecip_d", "llvm.loongarch.lasx.xvfrecip.s" => "__builtin_lasx_xvfrecip_s", + "llvm.loongarch.lasx.xvfrecipe.d" => "__builtin_lasx_xvfrecipe_d", + "llvm.loongarch.lasx.xvfrecipe.s" => "__builtin_lasx_xvfrecipe_s", "llvm.loongarch.lasx.xvfrint.d" => "__builtin_lasx_xvfrint_d", "llvm.loongarch.lasx.xvfrint.s" => "__builtin_lasx_xvfrint_s", "llvm.loongarch.lasx.xvfrintrm.d" => "__builtin_lasx_xvfrintrm_d", @@ -2541,6 +2551,8 @@ match name { "llvm.loongarch.lasx.xvfrintrz.s" => "__builtin_lasx_xvfrintrz_s", "llvm.loongarch.lasx.xvfrsqrt.d" => "__builtin_lasx_xvfrsqrt_d", "llvm.loongarch.lasx.xvfrsqrt.s" => "__builtin_lasx_xvfrsqrt_s", + "llvm.loongarch.lasx.xvfrsqrte.d" => "__builtin_lasx_xvfrsqrte_d", + "llvm.loongarch.lasx.xvfrsqrte.s" => "__builtin_lasx_xvfrsqrte_s", "llvm.loongarch.lasx.xvfrstp.b" => "__builtin_lasx_xvfrstp_b", "llvm.loongarch.lasx.xvfrstp.h" => "__builtin_lasx_xvfrstp_h", "llvm.loongarch.lasx.xvfrstpi.b" => "__builtin_lasx_xvfrstpi_b", @@ -3255,6 +3267,8 @@ match name { "llvm.loongarch.lsx.vfnmsub.s" => "__builtin_lsx_vfnmsub_s", "llvm.loongarch.lsx.vfrecip.d" => "__builtin_lsx_vfrecip_d", "llvm.loongarch.lsx.vfrecip.s" => "__builtin_lsx_vfrecip_s", + "llvm.loongarch.lsx.vfrecipe.d" => "__builtin_lsx_vfrecipe_d", + "llvm.loongarch.lsx.vfrecipe.s" => "__builtin_lsx_vfrecipe_s", "llvm.loongarch.lsx.vfrint.d" => "__builtin_lsx_vfrint_d", "llvm.loongarch.lsx.vfrint.s" => "__builtin_lsx_vfrint_s", "llvm.loongarch.lsx.vfrintrm.d" => "__builtin_lsx_vfrintrm_d", @@ -3267,6 +3281,8 @@ match name { "llvm.loongarch.lsx.vfrintrz.s" => "__builtin_lsx_vfrintrz_s", "llvm.loongarch.lsx.vfrsqrt.d" => "__builtin_lsx_vfrsqrt_d", "llvm.loongarch.lsx.vfrsqrt.s" => "__builtin_lsx_vfrsqrt_s", + "llvm.loongarch.lsx.vfrsqrte.d" => "__builtin_lsx_vfrsqrte_d", + "llvm.loongarch.lsx.vfrsqrte.s" => "__builtin_lsx_vfrsqrte_s", "llvm.loongarch.lsx.vfrstp.b" => "__builtin_lsx_vfrstp_b", "llvm.loongarch.lsx.vfrstp.h" => "__builtin_lsx_vfrstp_h", "llvm.loongarch.lsx.vfrstpi.b" => "__builtin_lsx_vfrstpi_b", @@ -4434,6 +4450,7 @@ match name { "llvm.nvvm.abs.bf16x2" => "__nvvm_abs_bf16x2", "llvm.nvvm.abs.i" => "__nvvm_abs_i", "llvm.nvvm.abs.ll" => "__nvvm_abs_ll", + "llvm.nvvm.activemask" => "__nvvm_activemask", "llvm.nvvm.add.rm.d" => "__nvvm_add_rm_d", "llvm.nvvm.add.rm.f" => "__nvvm_add_rm_f", "llvm.nvvm.add.rm.ftz.f" => "__nvvm_add_rm_ftz_f", @@ -4522,6 +4539,7 @@ match name { "llvm.nvvm.ex2.approx.d" => "__nvvm_ex2_approx_d", "llvm.nvvm.ex2.approx.f" => "__nvvm_ex2_approx_f", "llvm.nvvm.ex2.approx.ftz.f" => "__nvvm_ex2_approx_ftz_f", + "llvm.nvvm.exit" => "__nvvm_exit", "llvm.nvvm.f2bf16.rn" => "__nvvm_f2bf16_rn", "llvm.nvvm.f2bf16.rn.relu" => "__nvvm_f2bf16_rn_relu", "llvm.nvvm.f2bf16.rz" => "__nvvm_f2bf16_rz", @@ -4722,8 +4740,11 @@ match name { "llvm.nvvm.mul24.ui" => "__nvvm_mul24_ui", "llvm.nvvm.mulhi.i" => "__nvvm_mulhi_i", "llvm.nvvm.mulhi.ll" => "__nvvm_mulhi_ll", + "llvm.nvvm.mulhi.s" => "__nvvm_mulhi_s", "llvm.nvvm.mulhi.ui" => "__nvvm_mulhi_ui", "llvm.nvvm.mulhi.ull" => "__nvvm_mulhi_ull", + "llvm.nvvm.mulhi.us" => "__nvvm_mulhi_us", + "llvm.nvvm.nanosleep" => "__nvvm_nanosleep", "llvm.nvvm.neg.bf16" => "__nvvm_neg_bf16", "llvm.nvvm.neg.bf16x2" => "__nvvm_neg_bf16x2", "llvm.nvvm.popc.i" => "__nvvm_popc_i", @@ -4783,6 +4804,7 @@ match name { "llvm.nvvm.read.ptx.sreg.envreg7" => "__nvvm_read_ptx_sreg_envreg7", "llvm.nvvm.read.ptx.sreg.envreg8" => "__nvvm_read_ptx_sreg_envreg8", "llvm.nvvm.read.ptx.sreg.envreg9" => "__nvvm_read_ptx_sreg_envreg9", + "llvm.nvvm.read.ptx.sreg.globaltimer" => "__nvvm_read_ptx_sreg_globaltimer", "llvm.nvvm.read.ptx.sreg.gridid" => "__nvvm_read_ptx_sreg_gridid", // [DUPLICATE]: "llvm.nvvm.read.ptx.sreg.gridid" => "__nvvm_read_ptx_sreg_", "llvm.nvvm.read.ptx.sreg.laneid" => "__nvvm_read_ptx_sreg_laneid", @@ -4835,6 +4857,7 @@ match name { "llvm.nvvm.redux.sync.umax" => "__nvvm_redux_sync_umax", "llvm.nvvm.redux.sync.umin" => "__nvvm_redux_sync_umin", "llvm.nvvm.redux.sync.xor" => "__nvvm_redux_sync_xor", + "llvm.nvvm.reflect" => "__nvvm_reflect", "llvm.nvvm.rotate.b32" => "__nvvm_rotate_b32", "llvm.nvvm.rotate.b64" => "__nvvm_rotate_b64", "llvm.nvvm.rotate.right.b64" => "__nvvm_rotate_right_b64", @@ -4845,7 +4868,11 @@ match name { "llvm.nvvm.rsqrt.approx.f" => "__nvvm_rsqrt_approx_f", "llvm.nvvm.rsqrt.approx.ftz.f" => "__nvvm_rsqrt_approx_ftz_f", "llvm.nvvm.sad.i" => "__nvvm_sad_i", + "llvm.nvvm.sad.ll" => "__nvvm_sad_ll", + "llvm.nvvm.sad.s" => "__nvvm_sad_s", "llvm.nvvm.sad.ui" => "__nvvm_sad_ui", + "llvm.nvvm.sad.ull" => "__nvvm_sad_ull", + "llvm.nvvm.sad.us" => "__nvvm_sad_us", "llvm.nvvm.saturate.d" => "__nvvm_saturate_d", "llvm.nvvm.saturate.f" => "__nvvm_saturate_f", "llvm.nvvm.saturate.ftz.f" => "__nvvm_saturate_ftz_f", @@ -5471,6 +5498,7 @@ match name { "llvm.ppc.fctiwz" => "__builtin_ppc_fctiwz", "llvm.ppc.fctudz" => "__builtin_ppc_fctudz", "llvm.ppc.fctuwz" => "__builtin_ppc_fctuwz", + "llvm.ppc.fence" => "__builtin_ppc_fence", "llvm.ppc.fmaf128.round.to.odd" => "__builtin_fmaf128_round_to_odd", "llvm.ppc.fmsub" => "__builtin_ppc_fmsub", "llvm.ppc.fmsubs" => "__builtin_ppc_fmsubs", @@ -5599,6 +5627,9 @@ match name { "llvm.ppc.qpx.qvstfs" => "__builtin_qpx_qvstfs", "llvm.ppc.qpx.qvstfsa" => "__builtin_qpx_qvstfsa", "llvm.ppc.readflm" => "__builtin_readflm", + "llvm.ppc.rldimi" => "__builtin_ppc_rldimi", + "llvm.ppc.rlwimi" => "__builtin_ppc_rlwimi", + "llvm.ppc.rlwnm" => "__builtin_ppc_rlwnm", "llvm.ppc.scalar.extract.expq" => "__builtin_vsx_scalar_extract_expq", "llvm.ppc.scalar.insert.exp.qp" => "__builtin_vsx_scalar_insert_exp_qp", "llvm.ppc.set.texasr" => "__builtin_set_texasr", @@ -5912,6 +5943,8 @@ match name { "llvm.s390.vupllb" => "__builtin_s390_vupllb", "llvm.s390.vupllf" => "__builtin_s390_vupllf", "llvm.s390.vupllh" => "__builtin_s390_vupllh", + // spv + "llvm.spv.create.handle" => "__builtin_hlsl_create_handle", // ve "llvm.ve.vl.andm.MMM" => "__builtin_ve_vl_andm_MMM", "llvm.ve.vl.andm.mmm" => "__builtin_ve_vl_andm_mmm",