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
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 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()
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()
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 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()
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':
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 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()
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 execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == 'i32':
sdst = instruction[1]
simm16 = instruction[2]
if flag_of_status == OperationStatus.to_fill_node:
node.state.registers[sdst] = Register(simm16, Type.unknown,
Integrity.integer)
make_version(node.state, decompiler_data.versions, sdst)
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()
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 execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == 'i32':
sdst = instruction[1]
simm16 = instruction[2][instruction[2].find("x") + 1:]
if flag_of_status == OperationStatus.to_fill_node:
new_val, sdst_flag, simm16_flag = make_op(
node, sdst, simm16, " * ", '', '')
node.state.registers[sdst] = Register(new_val, Type.unknown,
Integrity.integer)
make_version(node.state, decompiler_data.versions, sdst)
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 == 'i32':
ssrc0 = instruction[1]
ssrc1 = instruction[2]
if flag_of_status == OperationStatus.to_fill_node:
node.state.registers["scc"] = \
Register('(int)' + node.state.registers[ssrc0].val + " < (int)" + node.state.registers[ssrc1].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 == "f32":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
new_val, src0_reg, src1_reg = make_op(node, src0, src1, " * ", 'as_float(', 'as_float(')
end_val = new_val + " + as_float(" + node.state.registers[vdst].val + ')'
node.state.registers[vdst] = Register(end_val, Type.unknown, Integrity.integer)
make_version(node.state, decompiler_data.versions, vdst)
node.state.registers[vdst].make_prev()
node.state.registers[vdst].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 == 'b32' or suffix == 'b64':
sdst = instruction[1]
ssrc0 = instruction[2]
if flag_of_status == OperationStatus.to_fill_node:
if node.state.registers.get(ssrc0) is not None:
node.state.registers[sdst] = \
Register(node.state.registers[ssrc0].val, node.state.registers[ssrc0].type, Integrity.integer)
else:
node.state.registers[sdst] = Register(ssrc0, Type.int32, Integrity.integer)
make_version(node.state, decompiler_data.versions, sdst)
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":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
new_integrity = node.state.registers[src1].integrity
new_val, src0_reg, src1_reg = make_op(node, src0, src1, " * ", '', '')
node.state.registers[vdst] = Register(new_val, Type.unknown, new_integrity)
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 execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
if suffix == "u32":
addr = instruction[1]
vdata0 = 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[name].val = \
node.state.registers[name].val + " + " + node.state.registers[vdata0].val
make_version(node.state, decompiler_data.versions, name)
node.state.registers[name].type_of_data = suffix
return node
output_string = name + " += " + node.state.registers[vdata0].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 = "min((int)" + node.state.registers[ssrc0].val + ", (int)" \
+ node.state.registers[ssrc1].val + ")"
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()
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_flag, ssrc1_flag = make_op(
node, ssrc0, str(pow(2, int(ssrc1))), " / ", '(int)', '')
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()
output_string = ""
if suffix == 'b64':
sdst = instruction[1]
ssrc0 = instruction[2]
ssrc1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
if ssrc0 == "exec":
ssrc0 = "1"
new_val = node.state.registers[
"scc"].val + " ? " + ssrc0 + " : " + ssrc1
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()
output_string = ""
if suffix == "u32":
vdst = instruction[1]
vcc = instruction[2]
src0 = instruction[3]
src1 = instruction[4]
if flag_of_status == OperationStatus.to_fill_node:
new_val, src0_reg, src1_reg = make_op(node, src0, src1, " - ", '(ulong)', '')
type_reg = Type.int32
if src0_reg:
type_reg = node.state.registers[src0].integrity
elif src1_reg:
type_reg = node.state.registers[src1].integrity
node.state.registers[vdst] = Register(new_val, Type.unknown, type_reg)
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
if suffix == "f32":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
new_val, src0_reg, src1_reg = make_op(node, src0, src1, " - ", 'as_float(', 'as_float(')
type_reg = Type.int32
if src0_reg:
type_reg = node.state.registers[src0].integrity
elif src1_reg:
type_reg = node.state.registers[src1].integrity
node.state.registers[vdst] = Register(new_val, Type.unknown, type_reg)
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 execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == "u32" or suffix == "i32" or suffix == 'f32':
if suffix != 'f32':
new_as_type = '(' + make_type(suffix) + ')'
else:
new_as_type = 'as_float('
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, " == ", new_as_type, new_as_type)
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
def execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == "u64": # not very correct
sdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
first_to, last_to, name_of_to, name_of_from, first_from, last_from \
= find_first_last_num_to_from(src0, src1)
src0 = name_of_to + str(first_to)
src1 = name_of_from + str(first_from)
new_val, src0_flag, src1_flag = make_op(
node, src0, src1, " > ", '(ulong)', '(uint)')
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
# decompiler_data.output_file.write(sdst + " = (ulong)" + src0 + " > (uint)" + src1 + "\n")
elif 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
def execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == "i32":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
if flag_of_status == OperationStatus.to_fill_node:
node.state.registers[vdst] = \
Register(node.state.registers[src1].val, node.state.registers[src1].type, Integrity.integer)
# node.state.registers[vdst].version = node.parent[0].state.registers[vdst].version
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
elif suffix == "i64":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
first_to, last_to, name_of_to, name_of_from, first_from, last_from \
= find_first_last_num_to_from(vdst, src1)
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 node.state.registers[from_registers].val == "0":
node.state.registers[from_registers].val = \
node.state.registers[name_of_from + str(first_from + 1)].val
new_val, src1_flag, src0_flag = make_op(
node, from_registers, str(pow(2, int(src0))), " / ", '',
'(long)')
node.state.registers[to_registers] = \
Register(new_val, node.state.registers[from_registers].type, Integrity.low_part)
# node.state.registers[to_registers].version = \
# node.parent[0].state.registers[to_registers].version
make_version(node.state, decompiler_data.versions,
to_registers)
node.state.registers[to_registers].type_of_data = suffix
to_registers_1 = name_of_to + str(last_to)
node.state.registers[to_registers_1] = \
Register(new_val, node.state.registers[from_registers].type, Integrity.high_part)
# node.state.registers[to_registers_1].version = \
# node.parent[0].state.registers[to_registers_1].version
make_version(node.state, decompiler_data.versions,
to_registers_1)
node.state.registers[to_registers_1].type_of_data = suffix
return node
return output_string
def upload_usesetup(state, to_registers, offset, parent):
to_registers1 = ""
separation = to_registers.find(":")
if separation != -1:
to_registers1 = "s" + to_registers[separation + 1:-1]
to_registers = "s" + to_registers[2:separation]
if offset == "0x0":
state.registers[to_registers] = Register(to_registers,
Type.general_setup,
Integrity.integer)
make_version(state, parent, to_registers)
elif offset == "0x4":
state.registers["s0"] = Register("get_local_size(0)",
Type.local_size_x, Integrity.integer)
make_version(state, parent, "s0")
state.registers["s1"] = Register("get_local_size(2)",
Type.local_size_z, Integrity.integer)
make_version(state, parent, "s1")
state.registers["s2"] = Register("get_global_size(0)",
Type.global_size_x, Integrity.integer)
make_version(state, parent, "s2")
state.registers["s3"] = Register("get_global_size(1)",
Type.global_size_y, Integrity.integer)
make_version(state, parent, "s2")
elif offset == "0xc":
state.registers[to_registers] = Register("get_global_size(0)",
Type.global_size_x,
Integrity.integer)
make_version(state, parent, to_registers)
if to_registers1 != "":
state.registers[to_registers1] = Register("get_global_size(1)",
Type.global_size_y,
Integrity.integer)
make_version(state, parent, to_registers1)
elif offset == "0x10":
state.registers[to_registers] = Register("get_global_size(1)",
Type.global_size_y,
Integrity.integer)
make_version(state, parent, to_registers)
if to_registers1 != "":
state.registers[to_registers1] = Register("get_global_size(2)",
Type.global_size_z,
Integrity.integer)
make_version(state, parent, to_registers1)
elif offset == "0x14":
state.registers[to_registers] = Register("get_global_size(2)",
Type.global_size_z,
Integrity.integer)
make_version(state, parent, to_registers)
def upload(state, to_registers, from_registers, offset, kernel_params, parent):
first_to, last_to, name_of_to, name_of_from, first_from, _ \
= find_first_last_num_to_from(to_registers, from_registers)
from_registers = name_of_from + str(first_from)
to_registers = name_of_to + str(first_to)
if state.registers[from_registers].type == Type.arguments_pointer:
if offset == "0x0":
state.registers[to_registers] = Register("get_global_offset(0)",
Type.global_offset_x,
Integrity.integer)
make_version(state, parent, to_registers)
state.registers[name_of_to + str(first_to + 1)] = Register(
"get_global_offset(0)", Type.global_offset_x,
Integrity.integer)
make_version(state, parent, name_of_to + str(first_to + 1))
if last_to - first_to > 1:
state.registers[name_of_to + str(last_to - 1)] = \
Register("get_global_offset(1)", Type.global_offset_y, Integrity.integer)
make_version(state, parent, name_of_to + str(last_to - 1))
state.registers[name_of_to + str(last_to)] = \
Register("get_global_offset(1)", Type.global_offset_y, Integrity.integer)
make_version(state, parent, name_of_to + str(last_to))
elif offset == "0x8":
state.registers[to_registers] = Register("get_global_offset(1)",
Type.global_offset_y,
Integrity.integer)
make_version(state, parent, to_registers)
state.registers[name_of_to + str(first_to + 1)] = \
Register("get_global_offset(1)", Type.global_offset_y, Integrity.integer)
make_version(state, parent, name_of_to + str(first_to + 1))
if last_to - first_to > 1:
state.registers[name_of_to + str(last_to - 1)] = \
Register("get_global_offset(2)", Type.global_offset_z, Integrity.integer)
make_version(state, parent, name_of_to + str(last_to - 1))
state.registers[name_of_to + str(last_to)] = \
Register("get_global_offset(2)", Type.global_offset_z, Integrity.integer)
make_version(state, parent, name_of_to + str(last_to))
elif offset == "0x10":
state.registers[to_registers] = Register("get_global_offset(2)",
Type.global_offset_z,
Integrity.integer)
make_version(state, parent, to_registers)
state.registers[name_of_to + str(first_to + 1)] = Register(
"get_global_offset(2)", Type.global_offset_z,
Integrity.integer)
make_version(state, parent, name_of_to + str(first_to + 1))
else:
for (reg, val) in kernel_params[offset]:
if val[0] == "*":
type_param = Type.paramA
val = val[1:]
else:
type_param = Type.param
state.registers[reg] = Register(val, type_param,
Integrity.integer)
make_version(state, parent, reg)
def execute(self, node, instruction, flag_of_status, suffix):
decompiler_data = DecompilerData()
output_string = ""
if suffix == "b32":
vdst = 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, src1, str(pow(2, int(src0))), " * ", '', '')
node.state.registers[vdst] = Register(
new_val, node.state.registers[src1].type,
Integrity.integer)
node.state.registers[vdst].version = node.parent[
0].state.registers[src1].version
node.state.registers[vdst].type_of_data = suffix
return node
return output_string
elif suffix == "b64":
vdst = instruction[1]
src0 = instruction[2]
src1 = instruction[3]
first_to, last_to, name_of_to, name_of_from, first_from, last_from \
= find_first_last_num_to_from(vdst, src1)
if "s" in name_of_from or "v" in name_of_from:
from_registers = name_of_from + str(first_from)
else:
from_registers = src1
to_registers = name_of_to + str(first_to)
if flag_of_status == OperationStatus.to_fill_node:
to_registers_1 = name_of_to + str(last_to)
from_registers_1 = name_of_from + str(first_from + 1)
new_val, src1_flag, src0_flag = make_op(
node, from_registers, str(pow(2, int(src0))), " * ", '',
'')
if src0_flag and src1_flag:
if node.state.registers[from_registers].type \
in [Type.global_id_x, Type.global_id_y, Type.global_id_z] \
and node.state.registers[name_of_to + str(first_from + 1)].val == "0":
node.state.registers[to_registers] = \
Register(new_val, node.state.registers[from_registers].type, Integrity.low_part)
make_version(node.state, decompiler_data.versions,
to_registers)
node.state.registers[
to_registers].type_of_data = suffix
node.state.registers[to_registers_1] = \
Register(node.state.registers[from_registers_1].val,
node.state.registers[from_registers].type,
Integrity.high_part)
make_version(node.state, decompiler_data.versions,
to_registers_1)
node.state.registers[
to_registers_1].type_of_data = suffix
else:
node.state.registers[
to_registers] = node.state.registers[
from_registers]
make_version(node.state, decompiler_data.versions,
to_registers)
node.state.registers[
to_registers_1] = node.state.registers[
from_registers_1]
make_version(node.state, decompiler_data.versions,
to_registers_1)
else:
type_reg = Type.int32
if src0_flag:
type_reg = node.state.registers[src0].type
if src1_flag:
type_reg = node.state.registers[from_registers].type
node.state.registers[to_registers] = Register(
new_val, type_reg, Integrity.low_part)
make_version(node.state, decompiler_data.versions,
to_registers)
node.state.registers[to_registers_1] = Register(
new_val, type_reg, Integrity.high_part)
make_version(node.state, decompiler_data.versions,
to_registers_1)
node.state.registers[to_registers].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":
vdst = instruction[1]
sdst = instruction[2]
src0 = instruction[3]
src1 = instruction[4]
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.work_group_id_x_local_size_offset and \
node.state.registers[src1].type == Type.work_item_id_x or \
node.state.registers[src0].type == Type.global_offset_x and \
node.state.registers[src1].type == Type.work_group_id_x_work_item_id:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = Register(
"get_global_id(0)", Type.global_id_x,
new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_y_local_size_offset and \
node.state.registers[src1].type == Type.work_item_id_y or \
node.state.registers[src0].type == Type.global_offset_y and \
node.state.registers[src1].type == Type.work_group_id_y_work_item_id:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = Register(
"get_global_id(1)", Type.global_id_y,
new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_z_local_size_offset and \
node.state.registers[src1].type == Type.work_item_id_z or \
node.state.registers[src0].type == Type.global_offset_z and \
node.state.registers[src1].type == Type.work_group_id_z_work_item_id or \
node.state.registers[src1].type == Type.global_offset_z and \
node.state.registers[src0].type == Type.work_group_id_z_work_item_id:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = Register(
"get_global_id(2)", Type.global_id_z,
new_integrity)
elif node.state.registers[src0].type == Type.paramA:
new_integrity = node.state.registers[src1].integrity
if decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "int" \
or decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "uint":
if "1073741824" in node.state.registers[src1].val:
new_value, src0_flag, src1_flag = make_op(
node, src1, "1073741824", " * ", '', '')
else:
new_value, src0_flag, src1_flag = make_op(
node, src1, "4", " / ", '', '')
new_val = node.state.registers[
src0].val + "[" + new_value + "]"
node.state.registers[vdst] = \
Register(new_val, Type.param_global_id_x, new_integrity)
elif decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "long" \
or decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "ulong":
new_value, src0_flag, src1_flag = make_op(
node, src1, "8", " / ", '', '')
new_val = node.state.registers[
src0].val + "[" + new_value + "]"
node.state.registers[vdst] = \
Register(new_val, Type.param_global_id_x, new_integrity)
elif decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "char" \
or decompiler_data.type_params.get("*" + node.state.registers[src0].val) == "uchar":
new_val = node.state.registers[
src0].val + "[" + node.state.registers[
src1].val + "]"
node.state.registers[vdst] = \
Register(new_val, Type.param_global_id_x, new_integrity)
elif decompiler_data.type_params.get(
"*" +
node.state.registers[src0].val) == "float":
new_value, src0_flag, src1_flag = make_op(
node, src1, "4", " / ", '', '')
new_val = node.state.registers[
src0].val + "[" + new_value + "]"
node.state.registers[vdst] = \
Register(new_val, Type.param_global_id_x, new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_x_local_size and \
node.state.registers[src1].type == Type.work_item_id_x:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = \
Register("get_global_id(0) - get_global_offset(0)", Type.work_group_id_x_work_item_id,
new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_y_local_size and \
node.state.registers[src1].type == Type.work_item_id_y:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = \
Register("get_global_id(1) - get_global_offset(1)", Type.work_group_id_y_work_item_id,
new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_z_local_size and \
node.state.registers[src1].type == Type.work_item_id_z:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = \
Register("get_global_id(2) - get_global_offset(2)", Type.work_group_id_z_work_item_id,
new_integrity)
elif node.state.registers[src0].type == Type.work_group_id_x_local_size and \
node.state.registers[src1].type == Type.work_item_id_x or \
node.state.registers[src1].type == Type.work_group_id_x_local_size and \
node.state.registers[src0].type == Type.work_item_id_x:
new_integrity = node.state.registers[src1].integrity
node.state.registers[vdst] = \
Register("get_global_id(0) - get_global_offset(0)", Type.unknown, 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
