Handle a syntax corner case where a def does not end with a ;

Author: GuillaumeGomez

src/intrinsic/archs.rs

@@ -110,18 +110,18 @@ match name {
     "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.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",
+    "llvm.amdgcn.s.incperflevel" => "__builtin_amdgcn_s_incperflevel",
     "llvm.amdgcn.s.memrealtime" => "__builtin_amdgcn_s_memrealtime",
     "llvm.amdgcn.s.memtime" => "__builtin_amdgcn_s_memtime",
     "llvm.amdgcn.s.sendmsg" => "__builtin_amdgcn_s_sendmsg",
     "llvm.amdgcn.s.sendmsghalt" => "__builtin_amdgcn_s_sendmsghalt",
     "llvm.amdgcn.s.setprio" => "__builtin_amdgcn_s_setprio",
     "llvm.amdgcn.s.setreg" => "__builtin_amdgcn_s_setreg",
-    "llvm.amdgcn.s.sleep" => "__builtin_amdgcn_s_decperflevel",
-    // [DUPLICATE]: "llvm.amdgcn.s.sleep" => "__builtin_amdgcn_s_incperflevel",
-    // [DUPLICATE]: "llvm.amdgcn.s.sleep" => "__builtin_amdgcn_s_sleep",
+    "llvm.amdgcn.s.sleep" => "__builtin_amdgcn_s_sleep",
     "llvm.amdgcn.s.waitcnt" => "__builtin_amdgcn_s_waitcnt",
     "llvm.amdgcn.sad.hi.u8" => "__builtin_amdgcn_sad_hi_u8",
     "llvm.amdgcn.sad.u16" => "__builtin_amdgcn_sad_u16",
@@ -2689,6 +2689,10 @@ match name {
     "llvm.ppc.vsx.xvresp" => "__builtin_vsx_xvresp",
     "llvm.ppc.vsx.xvrsqrtedp" => "__builtin_vsx_xvrsqrtedp",
     "llvm.ppc.vsx.xvrsqrtesp" => "__builtin_vsx_xvrsqrtesp",
+    "llvm.ppc.vsx.xxblendvb" => "__builtin_vsx_xxblendvb",
+    "llvm.ppc.vsx.xxblendvd" => "__builtin_vsx_xxblendvd",
+    "llvm.ppc.vsx.xxblendvh" => "__builtin_vsx_xxblendvh",
+    "llvm.ppc.vsx.xxblendvw" => "__builtin_vsx_xxblendvw",
     "llvm.ppc.vsx.xxpermx" => "__builtin_vsx_xxpermx",
     // ptx
     "llvm.ptx.bar.sync" => "__builtin_ptx_bar_sync",

tools/generate_intrinsics.py

@@ -60,13 +60,13 @@ def extract_instrinsics(intrinsics, file):
             current_arch = None
         elif line.startswith("def "):
             content = ""
-            while not content.endswith(";") and pos < len(lines):
+            while not content.endswith(";") and not content.endswith("}") and pos < len(lines):
                 line = lines[pos].split(" // ")[0].strip()
                 content += line
                 pos += 1
             entries = re.findall('GCCBuiltin<"(\\w+)">', content)
             if len(entries) > 0:
-                intrinsic = content.split(":")[0].split(" ")[1].strip()
+                intrinsic = content.split("def ")[1].strip().split(":")[0].strip()
                 intrinsic = intrinsic.split("_")
                 if len(intrinsic) < 2 or intrinsic[0] != "int":
                     continue