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
