def process_src(name_of_program, set_of_config, set_of_instructions): decompiler_data = DecompilerData() process_config(set_of_config, name_of_program) process_kernel_params(set_of_instructions) last_node = Node([""], decompiler_data.initial_state) curr_node = last_node last_node_state = decompiler_data.initial_state decompiler_data.cfg = last_node num = 0 while num < len(set_of_instructions): result_for_check = process_single_instruction(set_of_instructions, num, curr_node, last_node_state, last_node) if result_for_check is not None: num, curr_node, set_of_instructions, last_node, last_node_state = result_for_check else: return check_for_use_new_version() remove_unusable_versions() if decompiler_data.checked_variables != {} or decompiler_data.variables != {}: change_values() make_region_graph_from_cfg() process_region_graph() create_opencl_body()
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() if suffix == "b32": addr = instruction[1] vdata0 = instruction[2] offset = int(instruction[3][7:]) if len(instruction) == 4 else 0 if decompiler_data.lds_vars.get(offset) is None: decompiler_data.lds_vars[offset] = [ "lds" + str(decompiler_data.lds_var_number), "u" + suffix[1:] ] decompiler_data.lds_var_number += 1 new_value, src0_flag, src1_flag = make_op(node, addr, "4", " / ", '', '') name = decompiler_data.lds_vars[offset][0] + "[" + new_value + "]" if flag_of_status == OperationStatus.to_fill_node: node.state.registers[name] = \ Register(node.state.registers[vdata0].val, node.state.registers[vdata0].type, Integrity.integer) make_version(node.state, decompiler_data.versions, name) node.state.registers[name].type_of_data = "u" + suffix[1:] return node output_string = name + " = " + node.state.registers[name].val return output_string elif suffix == "b64": v = "v" + str(decompiler_data.number_of_v) addr = instruction[1] vdata0 = instruction[2] offset = instruction[3][7:] decompiler_data.output_file.write("ulong* " + v + "\n") decompiler_data.output_file.write(v + " = (ulong*)(ds + ((" + addr + " + " + offset + ") & ~3))\n") decompiler_data.output_file.write("*" + v + " = " + vdata0 + "\n") decompiler_data.number_of_v += 1
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() if suffix == "b32": vdst = instruction[1] src0 = instruction[2] src1 = instruction[3] ssrc2 = instruction[4] variable = "var" + str(decompiler_data.num_of_var) if flag_of_status == OperationStatus.to_fill_node: node.state.registers[vdst] = Register(variable, Type.program_param, Integrity.integer) make_version(node.state, decompiler_data.versions, vdst) if vdst in [src0, src1]: node.state.registers[vdst].make_prev() node.state.registers[vdst].type_of_data = suffix if node.state.registers[vdst].type == Type.param_global_id_x: variable = "*" + variable node.state.registers[vdst].val = variable decompiler_data.num_of_var += 1 decompiler_data.variables[ node.state.registers[vdst].version] = variable decompiler_data.names_of_vars[variable] = suffix return node if "?" in node.state.registers[ssrc2].val: node.state.registers[ ssrc2].val = "(" + node.state.registers[ssrc2].val + ")" output_string = node.state.registers[vdst].val + " = " + node.state.registers[ssrc2].val \ + " ? " + node.parent[0].state.registers[src1].val + " : " \ + node.parent[0].state.registers[src0].val return output_string
def get_kernel_params(offsets_of_kernel_params, offset_num): decompiler_data = DecompilerData() for key in sorted(offsets_of_kernel_params.keys()): instr = offsets_of_kernel_params[key][0] registers = offsets_of_kernel_params[key][1] if instr[-1] == 'd': num_output_regs = 1 first_num_of_reg = int(registers[1:]) else: num_output_regs = int(instr[-1]) first_num_of_reg = int(registers[2:registers.find(':')]) curr_reg = 0 name_of_reg = registers[0] offset_num_keys = sorted(offset_num.keys()) curr_num_offset = offset_num_keys.index(key) decompiler_data.kernel_params[key] = [] while curr_reg < num_output_regs: curr_offset = offset_num_keys[curr_num_offset] name_of_param = offset_num[curr_offset] decompiler_data.kernel_params[key].append( (name_of_reg + str(first_num_of_reg + curr_reg), name_of_param)) curr_reg += 1 if name_of_param[0] == "*" or ( decompiler_data.type_params.get(name_of_param) and 'long' in decompiler_data.type_params.get(name_of_param)): decompiler_data.kernel_params[key].append( (name_of_reg + str(first_num_of_reg + curr_reg), name_of_param)) curr_reg += 1 curr_num_offset += 1
def process_params(set_of_config, name_of_program): decompiler_data = DecompilerData() parameters = set_of_config[17:] for num_of_setting, set_of_config_num in enumerate(set_of_config): if ".arg" in set_of_config_num and "_." not in set_of_config_num: parameters = set_of_config[num_of_setting:] break decompiler_data.output_file.write("void " + name_of_program + "(") num_of_param = 0 flag_start = False for param in parameters: if not flag_start: flag_start = True else: decompiler_data.output_file.write(", ") set_of_param = param.strip().replace(',', ' ').split() name_param = set_of_param[1] type_param = set_of_param[3] flag_param = "" if len(set_of_param) > 4: flag_param = "__" + set_of_param[4] + " " if type_param[-1] == "*": name_param = "*" + name_param type_param = type_param[:-1] decompiler_data.params["param" + str(num_of_param)] = name_param decompiler_data.type_params[name_param] = type_param decompiler_data.output_file.write(flag_param + type_param + " " + name_param) num_of_param += 1 decompiler_data.output_file.write(")\n")
def update_reg_version(reg, curr_node, max_version, version_of_reg): decompiler_data = DecompilerData() if len(version_of_reg) == 1: if curr_node.state.registers.get( reg) is None or curr_node.state.registers[reg] is None: for parent in curr_node.parent: if parent.state.registers[reg] is not None: curr_node.state.registers[reg] = parent.state.registers[ reg] break elif len(version_of_reg) > 1: curr_node.state.registers[reg].version = reg + "_" + str(max_version + 1) curr_node.state.registers[reg].prev_version = list(version_of_reg) variable = "var" + str(decompiler_data.num_of_var) if curr_node.state.registers[reg].type == Type.param_global_id_x: variable = "*" + variable curr_node.state.registers[reg].val = variable decompiler_data.num_of_var += 1 for ver in version_of_reg: decompiler_data.variables[ver] = variable decompiler_data.checked_variables[ curr_node.state.registers[reg].version] = variable decompiler_data.versions[reg] = max_version + 1 return curr_node
def check_for_use_new_version_in_one_instruction(curr_node, instruction): decompiler_data = DecompilerData() for num_of_reg in range(1, len(curr_node.instruction)): register = curr_node.instruction[num_of_reg] if ("flat_store" in curr_node.instruction[0] or num_of_reg > 1) \ and len(register) > 1 \ and "cnd" not in curr_node.instruction[0]: if register[1] == "[": register = register[0] + register[2:register.find(":")] first_reg = curr_node.instruction[1] if first_reg[1] == "[": first_reg = first_reg[0] + first_reg[2:first_reg.find(":")] if curr_node.state.registers.get(register) is not None: checked_version = curr_node.state.registers[register].version else: continue if first_reg == register: checked_version = curr_node.parent[0].state.registers[ register].version if curr_node.state.registers.get(register) is not None \ and decompiler_data.checked_variables.get(checked_version) is not None \ and curr_node.state.registers[register].type_of_data is not None \ and (register != "vcc" or "and_saveexec" in instruction[0]): var_name = decompiler_data.checked_variables[checked_version] if decompiler_data.names_of_vars.get(var_name) is None: if decompiler_data.checked_variables.get( curr_node.parent[0].state.registers[register]. version) is not None: decompiler_data.names_of_vars[var_name] = \ curr_node.parent[0].state.registers[register].type_of_data else: decompiler_data.names_of_vars[var_name] = \ curr_node.state.registers[register].type_of_data
def process_size_of_work_groups(set_of_config): decompiler_data = DecompilerData() if ".cws" in set_of_config[1]: decompiler_data.size_of_work_groups = set_of_config[1].replace( ',', ' ').split()[1:] decompiler_data.output_file.write( "__kernel __attribute__((reqd_work_group_size(" + decompiler_data.size_of_work_groups[0] + ", " + decompiler_data.size_of_work_groups[1] + ", " + decompiler_data.size_of_work_groups[2] + ")))\n") else: decompiler_data.output_file.write("__kernel ")
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() reladdr = instruction[1] if decompiler_data.to_node.get(reladdr) is not None: node.add_child(decompiler_data.to_node[reladdr]) decompiler_data.to_node[reladdr].add_parent(node) else: if decompiler_data.from_node.get(reladdr) is None: decompiler_data.from_node[reladdr] = [node] else: decompiler_data.from_node[reladdr].append(node) node.instruction = "branch" return node
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() if suffix == "dword": vdst = instruction[1] vaddr = instruction[2] inst_offset = instruction[3] if len(instruction) > 3 else "" variable = "var" + str(decompiler_data.num_of_var) first_to, last_to, name_of_to, name_of_from, first_from, last_from \ = find_first_last_num_to_from(vdst, vaddr) from_registers = name_of_from + str(first_from) to_registers = name_of_to + str(first_to) if flag_of_status == OperationStatus.to_fill_node: if inst_offset == "": if first_to == last_to: data_type = node.state.registers[ from_registers].type_of_data node.state.registers[to_registers] = \ Register(variable, Type.program_param, Integrity.integer) make_version(node.state, decompiler_data.versions, to_registers) node.state.registers[ to_registers].type_of_data = data_type node.state.registers[to_registers].val = variable decompiler_data.num_of_var += 1 decompiler_data.variables[node.state.registers[ to_registers].version] = variable decompiler_data.names_of_vars[ variable] = node.state.registers[ from_registers].type_of_data return node output_string = node.state.registers[to_registers].val + " = " \ + node.parent[0].state.registers[from_registers].val return output_string elif suffix == "dwordx4": vdst = instruction[1] vaddr = instruction[2] inst_offset = instruction[3] vm = "vm" + str(decompiler_data.number_of_vm) decompiler_data.output_file.write("short* " + vm + " = (" + vaddr + " + " + inst_offset + ")\n") decompiler_data.output_file.write(vdst + "[0] = *(uint*)" + vm + "\n") decompiler_data.output_file.write(vdst + "[1] = *(uint*)(" + vm + " + 4)\n") decompiler_data.output_file.write(vdst + "[2] = *(uint*)(" + vm + " + 8)\n") decompiler_data.output_file.write(vdst + "[3] = *(uint*)(" + vm + " + 12)\n") decompiler_data.number_of_vm += 1
def make_region_graph_from_cfg(): decompiler_data = DecompilerData() curr_node = decompiler_data.cfg region = Region(TypeNode.linear, curr_node) decompiler_data.starts_regions[curr_node] = region decompiler_data.ends_regions[curr_node] = region visited = [curr_node] q = deque() q.append(curr_node.children[0]) while q: curr_node = q.popleft() if curr_node not in visited: visited.append(curr_node) if len(curr_node.parent) == 1 and (len(curr_node.children) == 1 or len(curr_node.children) == 0): if decompiler_data.ends_regions[ curr_node.parent[0]].type == TypeNode.linear: region = decompiler_data.ends_regions.pop( curr_node.parent[0]) region.end = curr_node decompiler_data.ends_regions[curr_node] = region else: region = Region(TypeNode.linear, curr_node) decompiler_data.starts_regions[curr_node] = region decompiler_data.ends_regions[curr_node] = region parent = decompiler_data.ends_regions[curr_node.parent[0]] parent.add_child(region) region.add_parent(parent) else: region = Region(TypeNode.basic, curr_node) decompiler_data.starts_regions[curr_node] = region decompiler_data.ends_regions[curr_node] = region flag_of_continue = False for c_p in curr_node.parent: if c_p not in visited: flag_of_continue = True break if flag_of_continue: visited.remove(curr_node) continue for c_p in curr_node.parent: if decompiler_data.ends_regions.get(c_p) is not None: parent = decompiler_data.ends_regions[c_p] else: parent = decompiler_data.ends_regions[ curr_node.parent[0].children[0]] parent.add_child(region) region.add_parent(parent) for child in curr_node.children: if child not in visited: q.append(child) else: region.add_child(decompiler_data.starts_regions[child]) decompiler_data.starts_regions[child].add_parent(region)
def change_cfg_for_else_structure(curr_node, instruction): decompiler_data = DecompilerData() for parents_of_label in decompiler_data.parents_of_label: parents_of_label.children.remove(decompiler_data.label) last_node = decompiler_data.parents_of_label[1] last_node_state = copy.deepcopy(decompiler_data.parents_of_label[1].state) decompiler_data.from_node[instruction[1]].remove(curr_node) from_node = instruction[1] if decompiler_data.from_node.get(from_node, None) is None: decompiler_data.from_node[from_node] = [] for parents_of_label in decompiler_data.parents_of_label: if parents_of_label != decompiler_data.parents_of_label[1]: decompiler_data.from_node[from_node].append(parents_of_label) decompiler_data.flag_of_else = False return last_node, last_node_state
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() reladdr = instruction[1] if flag_of_status == OperationStatus.to_fill_node: if decompiler_data.to_node.get(reladdr) is not None: node.add_child(decompiler_data.to_node[reladdr]) decompiler_data.to_node[reladdr].add_parent(node) else: if decompiler_data.from_node.get(reladdr) is None: decompiler_data.from_node[reladdr] = [node] else: decompiler_data.from_node[reladdr].append(node) return node output_string = node.state.registers["vcc"].val return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() if suffix == "b32": tab = " " tmp = "tmp" + str(decompiler_data.number_of_tmp) dst = instruction[1] addr = instruction[2] src = instruction[3] offset = instruction[4][7:] decompiler_data.output_file.write("uint64 " + tmp + "\n") decompiler_data.output_file.write( "for (short i = 0; i < 64; i++)\n") decompiler_data.output_file.write("{\n") decompiler_data.output_file.write(tab + "uint lane_id = " + addr + "[(i + (" + offset + " >> 2)) & 63]\n") decompiler_data.output_file.write( tab + tmp + "[i] = exec & (1ULL << lane_id) != 0) ? " + src + "[lane_id] : 0\n") decompiler_data.output_file.write("}\n") decompiler_data.output_file.write( "for (short i = 0; i < 64; i++)\n") decompiler_data.output_file.write(tab + "if (exec & (1ULL << i) != 0)\n") decompiler_data.output_file.write(tab + tab + dst + "[i] = " + tmp + "[i]\n") decompiler_data.number_of_tmp += 1
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() if suffix == 'b64': sdst = instruction[1] ssrc0 = instruction[2] decompiler_data.output_file.write(sdst + " = ~" + ssrc0 + "\n") decompiler_data.output_file.write("scc = " + sdst + " != 0\n")
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] if suffix == 'b32': if flag_of_status == OperationStatus.to_fill_node: if node.state.registers[ssrc0].type == Type.global_size_x \ and str(pow(2, int(ssrc1))) == decompiler_data.size_of_work_groups[0]: node.state.registers[sdst] = \ Register("get_num_groups(0)", node.state.registers[ssrc0].type, Integrity.integer) elif node.state.registers[ssrc0].type == Type.global_size_y \ and str(pow(2, int(ssrc1))) == decompiler_data.size_of_work_groups[1]: node.state.registers[sdst] = \ Register("get_num_groups(1)", node.state.registers[ssrc0].type, Integrity.integer) elif node.state.registers[ssrc0].type == Type.global_size_z \ and str(pow(2, int(ssrc1))) == decompiler_data.size_of_work_groups[2]: node.state.registers[sdst] = \ Register("get_num_groups(2)", node.state.registers[ssrc0].type, Integrity.integer) else: new_val, ssrc0_flag, ssrc1_flag = make_op( node, ssrc0, str(pow(2, int(ssrc1))), " / ", '', '') node.state.registers[sdst] = Register( new_val, node.state.registers[ssrc0].type, Integrity.integer) make_version(node.state, decompiler_data.versions, sdst) return node return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == "u32": vdst = instruction[1] sdst = instruction[2] src0 = instruction[3] src1 = instruction[4] ssrc2 = instruction[5] new_val, src0_reg, src1_reg = make_op(node, src0, src1, " + ", '(ulong)', '(ulong)') if flag_of_status == OperationStatus.to_fill_node: if src0_reg and src1_reg: if node.state.registers[src0].type == Type.paramA \ and node.state.registers[src1].type == Type.global_id_x: new_integrity = node.state.registers[src1].integrity node.state.registers[vdst] = \ Register(node.state.registers[src0].val + "[get_global_id(0)]", Type.param_global_id_x, new_integrity) else: new_integrity = node.state.registers[src1].integrity node.state.registers[vdst] = Register(new_val, Type.unknown, new_integrity) else: type_reg = Type.int32 if src0_reg: type_reg = node.state.registers[src0].type if src1_reg: type_reg = node.state.registers[src1].type node.state.registers[vdst] = Register(new_val, type_reg, Integrity.integer) make_version(node.state, decompiler_data.versions, vdst) if vdst in [src0, src1]: node.state.registers[vdst].make_prev() node.state.registers[vdst].type_of_data = suffix return node return output_string
def check_for_use_new_version(): decompiler_data = DecompilerData() curr_node = decompiler_data.cfg visited = [curr_node] queue = deque() queue.append(curr_node.children[0]) while queue: curr_node = queue.popleft() if curr_node not in visited: visited.append(curr_node) if len(curr_node.parent) < 2: if decompiler_data.checked_variables != {}: instruction = curr_node.instruction if instruction != "branch" and len(instruction) > 1: check_for_use_new_version_in_one_instruction( curr_node, instruction) else: flag_of_continue = False for c_p in curr_node.parent: if c_p not in visited: flag_of_continue = True break if flag_of_continue: visited.remove(curr_node) continue for child in curr_node.children: if child not in visited: queue.append(child)
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == 'i32' or suffix == 'u32': ssrc0 = instruction[1] ssrc1 = instruction[2] decompiler_data.output_file.write("scc = " + ssrc0 + " == " + ssrc1 + "\n") if suffix == 'u64': ssrc0 = instruction[1] ssrc1 = instruction[2] if flag_of_status == OperationStatus.to_fill_node: ssrc0 = ssrc0[0] + ssrc0[2:ssrc0.find(':')] if 's' in ssrc1: cmpr_val = node.state.registers[ssrc1].val else: cmpr_val = ssrc1 node.state.registers["scc"] = \ Register(node.state.registers[ssrc0].val + " == " + cmpr_val, Type.unknown, Integrity.integer) make_version(node.state, decompiler_data.versions, "scc") if "scc" in [ssrc0, ssrc1]: node.state.registers["scc"].make_prev() node.state.registers["scc"].type_of_data = suffix return node return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == "i32": sdst = instruction[1] src0 = instruction[2] src1 = instruction[3] if flag_of_status == OperationStatus.to_fill_node: new_val, src0_flag, src1_flag = make_op( node, src0, src1, " != ", '(int)', '(int)') node.state.registers[sdst] = Register(new_val, Type.unknown, Integrity.integer) make_version(node.state, decompiler_data.versions, sdst) if sdst in [src0, src1]: node.state.registers[sdst].make_prev() node.state.registers[sdst].type_of_data = suffix return node return output_string elif suffix == "u32": sdst = instruction[1] src0 = instruction[2] src1 = instruction[3] decompiler_data.output_file.write(sdst + " = (uint)" + src0 + " != (uint)" + src1 + "\n")
def change_values(): decompiler_data = DecompilerData() changes = {} curr_node = decompiler_data.cfg visited = [curr_node] queue = deque() queue.append(curr_node.children[0]) while queue: curr_node = queue.popleft() if curr_node not in visited: visited.append(curr_node) if len(curr_node.parent) < 2: instruction = curr_node.instruction if instruction != "branch" and len(instruction) > 1: change_values_for_one_instruction(curr_node, changes) else: flag_of_continue = False for c_p in curr_node.parent: if c_p not in visited: flag_of_continue = True break if flag_of_continue: visited.remove(curr_node) continue for child in curr_node.children: if child not in visited: queue.append(child)
def update_val_from_checked_variables(curr_node, register, check_version, first_reg, changes): decompiler_data = DecompilerData() instruction = curr_node.instruction if curr_node.state.registers.get(register) is not None \ and decompiler_data.variables.get(check_version) is not None \ and curr_node.state.registers[register].type_of_data is not None \ and (register != "vcc" or "and_saveexec" in instruction[0]): val_reg = curr_node.state.registers[register].val if register == first_reg: val_reg = curr_node.parent[0].state.registers[first_reg].val copy_val_prev = curr_node.state.registers[first_reg].val if decompiler_data.checked_variables.get(check_version) is not None: curr_node.state.registers[first_reg].val = \ curr_node.state.registers[first_reg].val.replace(val_reg, decompiler_data.checked_variables[check_version]) else: curr_node.state.registers[first_reg].val = \ curr_node.state.registers[first_reg].val.replace(val_reg, decompiler_data.variables[check_version]) copy_val_last = curr_node.state.registers[first_reg].val if copy_val_prev != copy_val_last: changes[curr_node.state.registers[first_reg].version] = [ copy_val_last, copy_val_prev ] update_value_for_reg(first_reg, curr_node)
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == 'i32': sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] temp = "temp" + str(decompiler_data.number_of_temp) decompiler_data.output_file.write(sdst + " = " + ssrc0 + " - " + ssrc1 + "\n") decompiler_data.output_file.write("long " + temp + " = (long)" + ssrc0 + " - (long)" + ssrc1 + "\n") decompiler_data.output_file.write("scc = " + temp + " > ((1LL << 31) - 1) || " + temp + " < (-1LL << 31)\n") decompiler_data.number_of_temp += 1 elif suffix == 'u32': sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] if flag_of_status == OperationStatus.to_fill_node: new_val, src0_reg, src1_reg = make_op(node, ssrc0, ssrc1, " - ", '(ulong)', '(ulong)') new_integrity = node.state.registers[ssrc1].integrity node.state.registers[sdst] = Register(new_val, Type.unknown, new_integrity) make_version(node.state, decompiler_data.versions, sdst) if sdst in [ssrc0, ssrc1]: node.state.registers[sdst].make_prev() node.state.registers[sdst].type_of_data = suffix return node return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == 'u32': sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] if flag_of_status == OperationStatus.to_fill_node: if ssrc1 == "0x20010": node.state.registers[sdst] = Register( "get_work_dim()", Type.work_dim, Integrity.integer) elif ssrc1 == '0x100010': node.state.registers[sdst] = Register( "get_local_size(1)", Type.local_size_y, Integrity.integer) else: print("Unknown pattern in s_bfe") make_version(node.state, decompiler_data.versions, sdst) return node return output_string if suffix == "i32": sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] if flag_of_status == OperationStatus.to_fill_node: return node return output_string
def get_offsets_to_regs(): decompiler_data = DecompilerData() offset_num = {} data_of_param = [] for num_of_param in range(len(decompiler_data.params)): num_of_param = str(num_of_param) name_of_param = decompiler_data.params["param" + num_of_param] type_of_param = decompiler_data.type_params.get(name_of_param) if type_of_param[-1] == '8': for i in range(8): data_of_param.append([ num_of_param, name_of_param + '.s' + str(i), type_of_param ]) else: data_of_param.append([num_of_param, name_of_param, type_of_param]) visited = False for num_of_param, name_of_param, type_of_param in data_of_param: if num_of_param == '0' and not visited: offset_num['0x30'] = name_of_param visited = True else: curr_offset = get_current_offset_for_not_first_param( offset_num, last_name, name_of_param, num_of_param) offset_num[curr_offset] = name_of_param last_name = name_of_param return offset_num
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == "b32": vdst = instruction[1] addr = instruction[2] offset = int(instruction[3][7:]) if len(instruction) == 4 else 0 new_value, src0_flag, src1_flag = make_op(node, addr, "4", " / ", '', '') name = decompiler_data.lds_vars[offset][0] + "[" + new_value + "]" if flag_of_status == OperationStatus.to_fill_node: node.state.registers[vdst] = Register(name, node.state.registers[name].type, Integrity.integer) make_version(node.state, decompiler_data.versions, vdst) node.state.registers[vdst].type_of_data = "u" + suffix[1:] return node return output_string elif suffix == "b64": vdst = instruction[1] addr = instruction[2] offset = int(instruction[3][7:]) if len(instruction) == 4 else 0 name = decompiler_data.lds_vars[offset][0] + "[" + node.state.registers[addr].var + "]" if flag_of_status == OperationStatus.to_fill_node: node.state.registers[vdst] = Register(name, node.state.registers[name].type, Integrity.integer) make_version(node.state, decompiler_data.versions, vdst) node.state.registers[vdst].type_of_data = "u" + suffix[1:] return node return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() vaddr = instruction[1] vdata = instruction[2] inst_offset = "0" if len(instruction) < 4 else instruction[3] first_to, last_to, name_of_to, name_of_from, first_from, last_from \ = find_first_last_num_to_from(vaddr, vdata) from_registers = name_of_from + str(first_from) to_registers = name_of_to + str(first_to) if flag_of_status == OperationStatus.to_fill_node: to_now = name_of_to + str(first_to + 1) node.state.registers[to_registers] = \ Register(node.state.registers[from_registers].val, node.state.registers[from_registers].type, Integrity.low_part) node.state.registers[to_registers].version = \ node.parent[0].state.registers[to_registers].version node.state.registers[to_registers].type_of_data = "dwordx2" second_from = name_of_from + str(first_from + 1) node.state.registers[to_now] = \ Register(node.state.registers[second_from].val, node.state.registers[second_from].type, Integrity.high_part) if node.parent[0].state.registers[to_now] is not None: node.state.registers[to_now].version = node.parent[ 0].state.registers[to_now].version node.state.registers[to_now].type_of_data = "dwordx2" return node else: var = node.parent[0].state.registers[to_registers].val if node.state.registers.get(from_registers): output_string = var + " = " + node.state.registers[ from_registers].val else: output_string = var + " = " + decompiler_data.initial_state.registers[ from_registers].val return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == 'i32': sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] if flag_of_status == OperationStatus.to_fill_node: new_val, ssrc0_reg, ssrc1_reg = make_op( node, ssrc0, ssrc1, " * ", '', '') if ssrc0_reg and ssrc1_reg: if node.state.registers[ssrc0].type == Type.local_size_x \ and node.state.registers[ssrc1].type == Type.work_group_id_x: node.state.registers[sdst] = \ Register(new_val, Type.work_group_id_x_local_size, Integrity.integer) elif node.state.registers[ssrc0].type == Type.local_size_y \ and node.state.registers[ssrc1].type == Type.work_group_id_y: node.state.registers[sdst] = \ Register(new_val, Type.work_group_id_y_local_size, Integrity.integer) elif node.state.registers[ssrc0].type == Type.local_size_z \ and node.state.registers[ssrc1].type == Type.work_group_id_z: node.state.registers[sdst] = \ Register(new_val, Type.work_group_id_z_local_size, Integrity.integer) else: node.state.registers[sdst] = Register( new_val, Type.unknown, Integrity.integer) else: node.state.registers[sdst] = Register( new_val, Type.unknown, Integrity.integer) make_version(node.state, decompiler_data.versions, sdst) if sdst in [ssrc0, ssrc1]: node.state.registers[sdst].make_prev() node.state.registers[sdst].type_of_data = suffix return node return output_string
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() vdst = instruction[1] src0 = instruction[2] src1 = instruction[3] decompiler_data.output_file.write(vdst + " = as_double(" + src0 + ") * as_double(" + src1 + ")\n")
def execute(self, node, instruction, flag_of_status, suffix): decompiler_data = DecompilerData() output_string = "" if suffix == 'u32': sdst = instruction[1] ssrc0 = instruction[2] ssrc1 = instruction[3] new_val, ssrc0_reg, ssrc1_reg = make_op(node, ssrc0, ssrc1, " + ", '(ulong)', '(ulong)') if flag_of_status == OperationStatus.to_fill_node: if ssrc0_reg and ssrc1_reg: node.state.registers[sdst] = \ Register(new_val, Type.unknown, Integrity.integer) else: type_reg = Type.int32 if ssrc0_reg: type_reg = node.state.registers[ssrc0].type if ssrc1_reg: type_reg = node.state.registers[ssrc1].type node.state.registers[sdst] = \ Register(new_val, type_reg, Integrity.integer) make_version(node.state, decompiler_data.versions, sdst) if sdst in [ssrc0, ssrc1]: node.state.registers[sdst].make_prev() node.state.registers[sdst].type_of_data = suffix return node return output_string