Merge pull request rust-lang#408 from GuillaumeGomez/intrinsics-conve…

…rsion

Update intrinsics conversion

src/intrinsic/archs.rs

@@ -151,20 +151,31 @@ match name {
     "llvm.amdgcn.msad.u8" => "__builtin_amdgcn_msad_u8",
     "llvm.amdgcn.perm" => "__builtin_amdgcn_perm",
     "llvm.amdgcn.permlane16" => "__builtin_amdgcn_permlane16",
+    "llvm.amdgcn.permlane16.var" => "__builtin_amdgcn_permlane16_var",
     "llvm.amdgcn.permlane64" => "__builtin_amdgcn_permlane64",
     "llvm.amdgcn.permlanex16" => "__builtin_amdgcn_permlanex16",
+    "llvm.amdgcn.permlanex16.var" => "__builtin_amdgcn_permlanex16_var",
     "llvm.amdgcn.qsad.pk.u16.u8" => "__builtin_amdgcn_qsad_pk_u16_u8",
     "llvm.amdgcn.queue.ptr" => "__builtin_amdgcn_queue_ptr",
     "llvm.amdgcn.rcp.legacy" => "__builtin_amdgcn_rcp_legacy",
     "llvm.amdgcn.readfirstlane" => "__builtin_amdgcn_readfirstlane",
     "llvm.amdgcn.readlane" => "__builtin_amdgcn_readlane",
     "llvm.amdgcn.rsq.legacy" => "__builtin_amdgcn_rsq_legacy",
     "llvm.amdgcn.s.barrier" => "__builtin_amdgcn_s_barrier",
+    "llvm.amdgcn.s.barrier.init" => "__builtin_amdgcn_s_barrier_init",
+    "llvm.amdgcn.s.barrier.join" => "__builtin_amdgcn_s_barrier_join",
+    "llvm.amdgcn.s.barrier.leave" => "__builtin_amdgcn_s_barrier_leave",
+    "llvm.amdgcn.s.barrier.signal" => "__builtin_amdgcn_s_barrier_signal",
+    "llvm.amdgcn.s.barrier.signal.isfirst" => "__builtin_amdgcn_s_barrier_signal_isfirst",
+    "llvm.amdgcn.s.barrier.signal.isfirst.var" => "__builtin_amdgcn_s_barrier_signal_isfirst_var",
+    "llvm.amdgcn.s.barrier.signal.var" => "__builtin_amdgcn_s_barrier_signal_var",
+    "llvm.amdgcn.s.barrier.wait" => "__builtin_amdgcn_s_barrier_wait",
     "llvm.amdgcn.s.dcache.inv" => "__builtin_amdgcn_s_dcache_inv",
     "llvm.amdgcn.s.dcache.inv.vol" => "__builtin_amdgcn_s_dcache_inv_vol",
     "llvm.amdgcn.s.dcache.wb" => "__builtin_amdgcn_s_dcache_wb",
     "llvm.amdgcn.s.dcache.wb.vol" => "__builtin_amdgcn_s_dcache_wb_vol",
     "llvm.amdgcn.s.decperflevel" => "__builtin_amdgcn_s_decperflevel",
+    "llvm.amdgcn.s.get.barrier.state" => "__builtin_amdgcn_s_get_barrier_state",
     "llvm.amdgcn.s.get.waveid.in.workgroup" => "__builtin_amdgcn_s_get_waveid_in_workgroup",
     "llvm.amdgcn.s.getpc" => "__builtin_amdgcn_s_getpc",
     "llvm.amdgcn.s.getreg" => "__builtin_amdgcn_s_getreg",
@@ -176,8 +187,10 @@ match name {
     "llvm.amdgcn.s.setprio" => "__builtin_amdgcn_s_setprio",
     "llvm.amdgcn.s.setreg" => "__builtin_amdgcn_s_setreg",
     "llvm.amdgcn.s.sleep" => "__builtin_amdgcn_s_sleep",
+    "llvm.amdgcn.s.sleep.var" => "__builtin_amdgcn_s_sleep_var",
     "llvm.amdgcn.s.wait.event.export.ready" => "__builtin_amdgcn_s_wait_event_export_ready",
     "llvm.amdgcn.s.waitcnt" => "__builtin_amdgcn_s_waitcnt",
+    "llvm.amdgcn.s.wakeup.barrier" => "__builtin_amdgcn_s_wakeup_barrier",
     "llvm.amdgcn.sad.hi.u8" => "__builtin_amdgcn_sad_hi_u8",
     "llvm.amdgcn.sad.u16" => "__builtin_amdgcn_sad_u16",
     "llvm.amdgcn.sad.u8" => "__builtin_amdgcn_sad_u8",
@@ -314,6 +327,8 @@ match name {
     // bpf
     "llvm.bpf.btf.type.id" => "__builtin_bpf_btf_type_id",
     "llvm.bpf.compare" => "__builtin_bpf_compare",
+    "llvm.bpf.getelementptr.and.load" => "__builtin_bpf_getelementptr_and_load",
+    "llvm.bpf.getelementptr.and.store" => "__builtin_bpf_getelementptr_and_store",
     "llvm.bpf.load.byte" => "__builtin_bpf_load_byte",
     "llvm.bpf.load.half" => "__builtin_bpf_load_half",
     "llvm.bpf.load.word" => "__builtin_bpf_load_word",
@@ -5776,14 +5791,6 @@ match name {
     "llvm.s390.verimf" => "__builtin_s390_verimf",
     "llvm.s390.verimg" => "__builtin_s390_verimg",
     "llvm.s390.verimh" => "__builtin_s390_verimh",
-    "llvm.s390.verllb" => "__builtin_s390_verllb",
-    "llvm.s390.verllf" => "__builtin_s390_verllf",
-    "llvm.s390.verllg" => "__builtin_s390_verllg",
-    "llvm.s390.verllh" => "__builtin_s390_verllh",
-    "llvm.s390.verllvb" => "__builtin_s390_verllvb",
-    "llvm.s390.verllvf" => "__builtin_s390_verllvf",
-    "llvm.s390.verllvg" => "__builtin_s390_verllvg",
-    "llvm.s390.verllvh" => "__builtin_s390_verllvh",
     "llvm.s390.vfaeb" => "__builtin_s390_vfaeb",
     "llvm.s390.vfaef" => "__builtin_s390_vfaef",
     "llvm.s390.vfaeh" => "__builtin_s390_vfaeh",
@@ -5815,7 +5822,7 @@ match name {
     "llvm.s390.vistrh" => "__builtin_s390_vistrh",
     "llvm.s390.vlbb" => "__builtin_s390_vlbb",
     "llvm.s390.vll" => "__builtin_s390_vll",
-    "llvm.s390.vlrl" => "__builtin_s390_vlrl",
+    "llvm.s390.vlrl" => "__builtin_s390_vlrlr",
     "llvm.s390.vmaeb" => "__builtin_s390_vmaeb",
     "llvm.s390.vmaef" => "__builtin_s390_vmaef",
     "llvm.s390.vmaeh" => "__builtin_s390_vmaeh",
@@ -5885,7 +5892,7 @@ match name {
     "llvm.s390.vstrczb" => "__builtin_s390_vstrczb",
     "llvm.s390.vstrczf" => "__builtin_s390_vstrczf",
     "llvm.s390.vstrczh" => "__builtin_s390_vstrczh",
-    "llvm.s390.vstrl" => "__builtin_s390_vstrl",
+    "llvm.s390.vstrl" => "__builtin_s390_vstrlr",
     "llvm.s390.vsumb" => "__builtin_s390_vsumb",
     "llvm.s390.vsumgf" => "__builtin_s390_vsumgf",
     "llvm.s390.vsumgh" => "__builtin_s390_vsumgh",