def set_init(self):
bar_spec = ptxILA.barSpec()
self.a_is_bar = (self.opcode_a == bar_spec.BAR_OPCODE)
#self.a_is_bar = (self.pc_a == 20)
self.a_previous_bar = self.model.bit('a_previous_bar')
self.model.set_next('a_previous_bar', self.a_is_bar)
self.a_cross_bar_flag = self.model.bit('a_cross_bar_flag')
self.model.set_next(
'a_cross_bar_flag',
ila.ite((self.a_is_bar == ila.bool(False)) &
(self.a_previous_bar == ila.bool(True)), ila.bool(True),
self.a_cross_bar_flag))
self.b_is_bar = (self.opcode_b == bar_spec.BAR_OPCODE)
self.b_come_to_bar_flag = self.model.bit('b_come_to_bar_flag')
self.model.set_next(
'b_come_to_bar_flag',
ila.ite(self.b_is_bar, ila.bool(True), self.b_come_to_bar_flag))
self.predicate_one = self.model.bit('predicate_one')
self.model.set_next('predicate_one', (~self.a_cross_bar_flag) |
(self.b_come_to_bar_flag))
self.model.set_init('a_previous_bar', self.model.bool(False))
self.model.set_init('a_cross_bar_flag', self.model.bool(False))
self.model.set_init('b_come_to_bar_flag', self.model.bool(False))
self.model.set_init('predicate_one', self.model.bool(True))
golden = ila.Abstraction('golden')
g_prop = golden.bit('prop')
golden.set_next('prop', golden.bool(True))
golden.set_init('prop', golden.bool(True))
ila.bmc(30, self.model, self.predicate_one, 1, golden, g_prop)
def pc_nxt(self):
bar_spec = ptxILA.barSpec()
self.pc_a_next = ila.ite(
self.arb == self.arbA,
ila.ite(
self.opcode_a == bar_spec.BAR_OPCODE,
ila.ite(self.bar_state_a == bar_spec.BAR_FINISH, self.pc_a + 4,
self.pc_a), self.pc_a + 4), self.pc_a)
self.pc_b_next = ila.ite(
self.arb == self.arbB,
ila.ite(
self.opcode_b == bar_spec.BAR_OPCODE,
ila.ite(self.bar_state_b == bar_spec.BAR_FINISH, self.pc_b + 4,
self.pc_b), self.pc_b + 4), self.pc_b)
self.model.set_next('pc_a', self.pc_a_next)
self.model.set_next('pc_b', self.pc_b_next)
def property_test2(self):
bar_spec = ptxILA.barSpec()
self.a_is_second_bar = (self.pc_a == 12 * instruction_format.MEM_BITS /
8)
self.a_previous_is_second_bar = self.model.bit(
'a_previous_is_second_bar')
self.model.set_next('a_previous_is_second_bar', self.a_is_second_bar)
self.a_cross_second_bar_flag = self.model.bit(
'a_cross_second_bar_flag')
self.model.set_next(
'a_cross_second_bar_flag',
ila.ite((self.a_is_second_bar == ila.bool(False)) &
(self.a_previous_is_second_bar == ila.bool(True)),
ila.bool(True), self.a_cross_second_bar_flag))
self.b_is_second_bar = (self.pc_b == ila.const(
12 * instruction_format.MEM_BITS / 8, instruction_format.PC_BITS))
self.b_come_to_second_bar_flag = self.model.bit(
'b_come_to_second_bar_flag')
self.model.set_next(
'b_come_to_second_bar_flag',
ila.ite(self.b_is_second_bar, ila.bool(True),
self.b_come_to_second_bar_flag))
self.predicate_two = self.model.bit('predicate_two')
self.model.set_next('predicate_two', (~self.a_cross_second_bar_flag) |
(self.b_come_to_second_bar_flag))
self.model.set_init('a_previous_is_second_bar', self.model.bool(False))
self.model.set_init('a_cross_second_bar_flag', self.model.bool(False))
self.model.set_init('b_come_to_second_bar_flag',
self.model.bool(False))
self.model.set_init('predicate_two', self.model.bool(True))
golden = ila.Abstraction('golden')
g_prop = golden.bit('prop')
golden.set_next('prop', golden.bool(True))
golden.set_init('prop', golden.bool(True))
ila.bmc(10, self.model, self.predicate_two, 1, golden, g_prop)
def createBar(self):
bar_spec = ptxILA.barSpec()
self.bar_state_a = self.model.reg('bar_state_a',
bar_spec.BAR_STATE_BITS)
self.bar_state_b = self.model.reg('bar_state_b',
bar_spec.BAR_STATE_BITS)
self.bar_counter_enter = self.model.reg(
'bar_counter_enter', bar_spec.BAR_COUNTER_ENTER_BITS)
self.bar_counter_exit = self.model.reg('bar_counter_exit',
bar_spec.BAR_COUNTER_EXIT_BITS)
self.bar_counter_max = bar_spec.THREAD_NUM
self.bar_state_a_next = ila.ite(self.arb == self.arbA, ila.ite(self.opcode_a == bar_spec.BAR_OPCODE, ila.ite(self.bar_state_a == bar_spec.BAR_FINISH, ila.const(bar_spec.BAR_INIT, bar_spec.BAR_STATE_BITS),\
ila.ite(self.bar_state_a == bar_spec.BAR_INIT, ila.const(bar_spec.BAR_ENTER, bar_spec.BAR_STATE_BITS),\
ila.ite(self.bar_state_a == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit != 0, self.bar_state_a,\
ila.ite(self.bar_counter_enter == (self.bar_counter_max - 1), ila.const(bar_spec.BAR_EXIT, bar_spec.BAR_STATE_BITS), ila.const(bar_spec.BAR_WAIT, bar_spec.BAR_STATE_BITS))),\
ila.ite(self.bar_state_a == bar_spec.BAR_WAIT, ila.ite(self.bar_counter_enter == self.bar_counter_max, ila.const(bar_spec.BAR_EXIT, bar_spec.BAR_STATE_BITS), self.bar_state_a),\
ila.ite(self.bar_state_a == bar_spec.BAR_EXIT, ila.const(bar_spec.BAR_FINISH, bar_spec.BAR_STATE_BITS), self.bar_state_a))))), self.bar_state_a), self.bar_state_a)
self.bar_state_b_next = ila.ite(self.arb == self.arbB, ila.ite(self.opcode_b == bar_spec.BAR_OPCODE, ila.ite(self.bar_state_b == bar_spec.BAR_FINISH, ila.const(bar_spec.BAR_INIT, bar_spec.BAR_STATE_BITS),\
ila.ite(self.bar_state_b == bar_spec.BAR_INIT, ila.const(bar_spec.BAR_ENTER, bar_spec.BAR_STATE_BITS),\
ila.ite(self.bar_state_b == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit != 0, self.bar_state_b,\
ila.ite(self.bar_counter_enter == (self.bar_counter_max - 1), ila.const(bar_spec.BAR_EXIT, bar_spec.BAR_STATE_BITS), ila.const(bar_spec.BAR_WAIT, bar_spec.BAR_STATE_BITS))),\
ila.ite(self.bar_state_b == bar_spec.BAR_WAIT, ila.ite(self.bar_counter_enter == self.bar_counter_max, ila.const(bar_spec.BAR_EXIT, bar_spec.BAR_STATE_BITS), self.bar_state_b),\
ila.ite(self.bar_state_b == bar_spec.BAR_EXIT, ila.const(bar_spec.BAR_FINISH, bar_spec.BAR_STATE_BITS), self.bar_state_b))))), self.bar_state_b), self.bar_state_b)
self.bar_counter_enter_next_a = ila.ite(self.opcode_a == bar_spec.BAR_OPCODE, \
ila.ite(self.bar_state_a == bar_spec.BAR_INIT, self.bar_counter_enter, \
ila.ite(self.bar_state_a == bar_spec.BAR_FINISH, self.bar_counter_enter,\
ila.ite(self.bar_state_a == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit == 0, self.bar_counter_enter + 1, self.bar_counter_enter),\
ila.ite(self.bar_state_a == bar_spec.BAR_WAIT, self.bar_counter_enter, \
ila.ite(self.bar_state_a == bar_spec.BAR_EXIT, ila.ite(self.bar_counter_exit == 1, ila.const(0x0, bar_spec.BAR_COUNTER_ENTER_BITS), self.bar_counter_enter), self.bar_counter_enter))))), \
self.bar_counter_enter)
self.bar_counter_enter_next_b = ila.ite(self.opcode_b == bar_spec.BAR_OPCODE, \
ila.ite(self.bar_state_b == bar_spec.BAR_INIT, self.bar_counter_enter, \
ila.ite(self.bar_state_b == bar_spec.BAR_FINISH, self.bar_counter_enter,\
ila.ite(self.bar_state_b == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit == 0, self.bar_counter_enter + 1, self.bar_counter_enter),\
ila.ite(self.bar_state_b == bar_spec.BAR_WAIT, self.bar_counter_enter, \
ila.ite(self.bar_state_b == bar_spec.BAR_EXIT, ila.ite(self.bar_counter_exit == 1, ila.const(0x0, bar_spec.BAR_COUNTER_ENTER_BITS), self.bar_counter_enter), self.bar_counter_enter))))), \
self.bar_counter_enter)
self.bar_counter_enter_next = ila.ite(self.arb == self.arbA,
self.bar_counter_enter_next_a,
self.bar_counter_enter_next_b)
self.bar_counter_exit_next_a = ila.ite(self.opcode_a == bar_spec.BAR_OPCODE,\
ila.ite(self.bar_state_a == bar_spec.BAR_INIT, self.bar_counter_exit,\
ila.ite(self.bar_state_a == bar_spec.BAR_FINISH, self.bar_counter_exit,\
ila.ite(self.bar_state_a == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit == 0, ila.ite(self.bar_counter_enter == (self.bar_counter_max - 1), ila.const(self.bar_counter_max, bar_spec.BAR_COUNTER_EXIT_BITS), self.bar_counter_exit), self.bar_counter_exit),\
ila.ite(self.bar_state_a == bar_spec.BAR_WAIT, self.bar_counter_exit,\
ila.ite(self.bar_state_a == bar_spec.BAR_EXIT, ila.ite(self.bar_counter_exit != 0, self.bar_counter_exit - 1, self.bar_counter_exit - 1), self.bar_counter_exit))))),\
self.bar_counter_exit)
self.bar_counter_exit_next_b = ila.ite(self.opcode_b == bar_spec.BAR_OPCODE,\
ila.ite(self.bar_state_b == bar_spec.BAR_INIT, self.bar_counter_exit,\
ila.ite(self.bar_state_b == bar_spec.BAR_FINISH, self.bar_counter_exit,\
ila.ite(self.bar_state_b == bar_spec.BAR_ENTER, ila.ite(self.bar_counter_exit == 0, ila.ite(self.bar_counter_enter == (self.bar_counter_max - 1), ila.const(self.bar_counter_max, bar_spec.BAR_COUNTER_EXIT_BITS), self.bar_counter_exit), self.bar_counter_exit),\
ila.ite(self.bar_state_b == bar_spec.BAR_WAIT, self.bar_counter_exit,\
ila.ite(self.bar_state_b == bar_spec.BAR_EXIT, ila.ite(self.bar_counter_exit != 0, self.bar_counter_exit - 1, self.bar_counter_exit - 1), self.bar_counter_exit))))),\
self.bar_counter_exit)
self.bar_counter_exit_next = ila.ite(self.arb == self.arbA,
self.bar_counter_exit_next_a,
self.bar_counter_exit_next_b)
self.model.set_next('bar_state_a', self.bar_state_a_next)
self.model.set_next('bar_state_b', self.bar_state_b_next)
self.model.set_next('bar_counter_enter', self.bar_counter_enter_next)
self.model.set_next('bar_counter_exit', self.bar_counter_exit_next)
self.model.set_init(
'bar_state_a',
self.model.const(bar_spec.BAR_INIT, bar_spec.BAR_STATE_BITS))
self.model.set_init(
'bar_state_b',
self.model.const(bar_spec.BAR_INIT, bar_spec.BAR_STATE_BITS))
self.model.set_init(
'bar_counter_enter',
self.model.const(0x0, bar_spec.BAR_COUNTER_ENTER_BITS))
self.model.set_init(
'bar_counter_exit',
self.model.const(0x0, bar_spec.BAR_COUNTER_EXIT_BITS))
def ptx_next_state(self, state):
mem = state['mem']
pc = state['pc']
instruction = mem[pc / 4]
print instruction
#pc += 4
#state['pc'] = pc
opcode = self.OPCODE_MASK & instruction
opcode = opcode >> instruction_format.OPCODE_BIT_BOT
dst = self.DST_MASK & instruction
dst = dst >> instruction_format.DST_BIT_BOT
src0 = self.SRC0_MASK & instruction
src0 = src0 >> instruction_format.SRC0_BIT_BOT
src1 = self.SRC1_MASK & instruction
src1 = src1 >> instruction_format.SRC1_BIT_BOT
base = self.BASE_MASK & instruction
base = base >> instruction_format.BASE_BIT_BOT
pred = self.P_REG_MASK & instruction
pred = pred >> self.P_REG_BIT
bra = (self.BRA_MASK & instruction) >> instruction_format.IMM_BIT_BOT
test_program = []
general_reg_book_file = 'general_reg_book'
general_reg_book_obj = open(general_reg_book_file)
general_reg_book = pickle.load(general_reg_book_obj)
for general_reg in general_reg_book:
test_program.append('mov.s32 ' + general_reg + ',' +
str(state[general_reg]) + '; ')
reg_book_file = 'reg_book'
reg_book_obj = open(reg_book_file, 'r')
reg_book = pickle.load(reg_book_obj)
instruction_book_file = 'instruction_book'
instruction_book_obj = open(instruction_book_file, 'r')
instruction_book = instruction_book_obj.readlines()
long_int_reg_book_file = 'long_int_reg_book'
long_int_reg_book_obj = open(long_int_reg_book_file, 'r')
long_int_reg_book = pickle.load(long_int_reg_book_obj)
long_int_reg_book_obj.close()
reg_book = general_reg_book + long_int_reg_book
#if ((opcode != self.OPCODE_MUL) & (opcode != self.OPCODE_ADD) & (opcode != self.OPCODE_SUB)):
if ((opcode != self.OPCODE_ADD) & (opcode != self.OPCODE_SUB) &
(opcode != self.OPCODE_BRA) & (opcode != self.OPCODE_BAR) &
(opcode != self.OPCODE_LD) & (opcode != self.OPCODE_ST) &
(opcode != self.OPCODE_MOV) & (opcode != self.OPCODE_MUL)):
state['pc'] = state['pc'] + 4
return state
if (opcode == self.OPCODE_BRA):
if base:
if pred >= len(reg_book):
return status
pred_reg_text = reg_book[pred]
if pred_reg_text not in general_reg_book:
return status
pred_reg_data = state[pred_reg_text]
if pred_reg_data:
pc += bra
state['pc'] = pc
else:
pc += bra
state['pc'] = pc
return state
'''
if (opcode == self.OPCODE_BAR):
bar_state = state['bar_state']
bar_spec = ptxILA.barSpec();
bar_counter_enter = state['bar_counter_enter']
bar_counter_exit = state['bar_counter_exit']
if (bar_state == bar_spec.BAR_FINISH):
state['pc'] = state['pc'] + 4
sim_program_line = ''
sim_program_line += 'mov.u32 %r1, ' + str(bar_state) + ';'
sim_program_line += 'mov.u32 %r23, ' + str(bar_counter_enter) + ';'
sim_program_line += 'mov.u32 %r24, ' + str(bar_counter_exit) + '; \n'
example_sim_program_file = 'tbar.ptx'
example_sim_program_obj = open(example_sim_program_file, 'r')
example_sim_program = example_sim_program_obj.readlines()
sim_program = []
bar_program_hole = 42
for i in range(len(example_sim_program)):
if i == bar_program_hole:
sim_program.append(sim_program_line)
else:
sim_program.append(example_sim_program[i])
example_sim_program_obj.close()
sim_program_file = 'tbar.ptx'
sim_program_obj = open(sim_program_file, 'w')
for sim_line in sim_program:
sim_program_obj.write(sim_line)
sim_program_obj.close()
(status, output) = commands.getstatusoutput('./dryrun_bar.out')
print status
print output
(status, output) = commands.getstatusoutput('sbatch parallel_bar.cmd')
print status
print output
output_word = output.split()
taskTag = output_word[3]
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' + taskTag + '.out')
while(status == 256):
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' + taskTag + '.out')
[bar_state, bar_counter_enter, bar_counter_exit] = output.split()
bar_state = int(bar_state)
bar_counter_enter = int(bar_counter_enter)
bar_counter_exit = int(bar_counter_exit)
if (bar_counter_enter < 0):
bar_counter_enter = -bar_counter_enter
bar_counter_enter = (1<<31) - bar_counter_enter + (1<<31)
if (bar_counter_exit < 0):
bar_counter_exit = -bar_counter_exit
bar_counter_exit = (1<<31) - bar_counter_exit + (1<<31)
state['bar_state'] = bar_state
state['bar_counter_enter'] = bar_counter_enter
state['bar_counter_exit'] = bar_counter_exit
return state
'''
if (opcode == self.OPCODE_BAR):
bar_state = state['bar_state']
#bar_counter_enter = state['bar_counter_enter']
#bar_counter_exit = state['bar_counter_exit']
bar_spec = ptxILA.barSpec()
bar_counter_enter = state['bar_counter_enter']
bar_counter_exit = state['bar_counter_exit']
if bar_state == bar_spec.BAR_INIT:
bar_state = bar_spec.BAR_ENTER
elif bar_state == bar_spec.BAR_FINISH:
state['pc'] = state['pc'] + 4
bar_state = bar_spec.BAR_INIT
elif bar_state == bar_spec.BAR_ENTER:
if (bar_counter_exit == 0):
bar_counter_enter = bar_counter_enter + 1
if bar_counter_enter == bar_spec.THREAD_NUM:
bar_state = bar_spec.BAR_EXIT
bar_counter_exit = bar_spec.THREAD_NUM
else:
if bar_counter_enter > bar_spec.THREAD_NUM:
state['bar_state'] = bar_spec.BAR_WAIT
return state
bar_state = bar_spec.BAR_WAIT
elif bar_state == bar_spec.BAR_WAIT:
if bar_counter_enter == bar_spec.THREAD_NUM:
bar_state = bar_spec.BAR_EXIT
elif bar_state == bar_spec.BAR_EXIT:
bar_counter_exit -= 1
bar_state = bar_spec.BAR_FINISH
if bar_counter_exit < 0:
state['bar_state'] = bar_spec.BAR_FINISH
if (bar_counter_exit < 0):
bar_counter_exit = -bar_counter_exit
bar_counter_exit = (
1 << (bar_spec.BAR_COUNTER_EXIT_BITS -
1)) - bar_counter_exit + (
1 << (bar_spec.BAR_COUNTER_EXIT - 1))
state['bar_counter_exit'] = bar_counter_exit
return state
if bar_counter_exit == 0:
bar_counter_enter = 0
state['bar_state'] = bar_state
state['bar_counter_enter'] = bar_counter_enter
state['bar_counter_exit'] = bar_counter_exit
return state
'''
if (bar_micro_flag):
bar_counter_enter = state['bar_counter_enter']
bar_counter_exit = state['bar_counter_exit']
if bar_state == bar_spec.BAR_ENTER:
if (bar_counter_exit == 0):
bar_counter_enter = bar_counter_enter + 1
if bar_counter_enter == bar_spec.THREAD_NUM:
bar_state = bar_spec.BAR_EXIT
bar_counter_exit = bar_spec.THREAD_NUM
else:
bar_state = bar_spec.BAR_WAIT
elif bar_state == bar_spec.BAR_WAIT:
if bar_counter_enter == bar_spec.THREAD_NUM:
bar_state = bar_spec.BAR_EXIT
elif bar_state == bar_spec.BAR_EXIT:
bar_counter_exit -= 1
bar_state = bar_spec.BAR_FINISH
if bar_counter_exit == 0:
bar_counter_enter = 0
state['bar_state'] = bar_state
state['bar_counter_enter'] = bar_counter_enter
state['bar_counter_exit'] = bar_counter_exit
else:
if bar_state == bar_spec.BAR_INIT:
bar_state = bar_spec.BAR_ENTER
elif bar_state == bar_spec.BAR_FINISH:
bar_state = bar_spec.BAR_INIT
state['pc'] = state['pc'] + 4
state['bar_state'] = bar_state
#state['bar_counter_enter'] = bar_counter_enter
#state['bar_counter_exit'] = bar_counter_exit
return state
'''
pc = pc + 4
state['pc'] = pc
op_text = instruction_book[opcode]
op_text = op_text[:(len(op_text) - 1)]
def find_addr(laddr):
for mem_key in mem_map.keys():
if (mem_map[mem_key][1]) >= laddr:
start_addr = mem_map[mem_key][0]
dmem_name = mem_key
return [dmem_name, start_addr]
if opcode == self.OPCODE_LD:
mem = state['dmem']
default = mem.default
values = mem.values
addr = []
value = []
for (a, v) in values:
addr.append(a * 4)
value.append(v)
dest_text = reg_book[dst]
self.ldAddr = (self.ldIMM_MASK
& instruction) >> instruction_format.IMM_BIT_BOT
self.ldAddr = (self.ldAddr) << 2
print 'load_addr' + str(self.ldAddr)
[dmem_name, start_addr] = find_addr(self.ldAddr)
item = (self.ldAddr - start_addr) >> 2
pre_ld_program = ''
pre_ld_program += '.reg .b64 %r_sim_ld<3>; .reg .b32 %r_ssim_ld;'
for i in range(len(addr)):
pre_addr = addr[i]
[pre_mem_name, pre_start_addr] = find_addr(pre_addr)
pre_item = (pre_addr - pre_start_addr) >> 2
pre_ld_program += 'ld.param.u64 %r_sim_ld1, [' + pre_mem_name + ']; '
pre_ld_program += 'cvta.to.global.u64 %r_sim_ld2, %r_sim_ld1; '
pre_ld_program += 'mov.u32 %r_ssim_ld, ' + str(pre_item) + '; '
pre_ld_program += 'mul.wide.s32 %r_sim_ld1, %r_ssim_ld, 4; '
pre_ld_program += 'add.s64 %r_sim_ld2, %r_sim_ld1, %r_sim_ld2; '
pre_ld_program += 'mov.u32 %r_ssim_ld, ' + str(value[i]) + '; '
pre_ld_program += 'st.global.b32 [%r_sim_ld2], %r_ssim_ld; '
pre_ld_program += '\n'
ld_program = ''
ld_program += 'ld.param.u64 %r_sim_ld1, [' + dmem_name + ']; '
ld_program += 'cvta.to.global.u64 %r_sim_ld2, %r_sim_ld1; '
ld_program += 'mov.u32 %r_ssim_ld, ' + str(item) + '; '
ld_program += 'mul.wide.s32 %r_sim_ld1, %r_ssim_ld, 4; '
ld_program += 'add.s64 %r_sim_ld2, %r_sim_ld1, %r_sim_ld2; '
ld_program += 'ld.global.b32 ' + dest_text + ',[%r_sim_ld2]; '
ld_program += 'mov.s32 %r9, ' + dest_text + '; '
example_sim_program_file = 't266.ptx'
example_sim_program_obj = open(example_sim_program_file, 'r')
example_sim_program = example_sim_program_obj.readlines()
sim_program = []
for test_program_line in test_program:
ld_program = test_program_line + ld_program
for i in range(len(example_sim_program)):
if i == self.EXAMPLE_PROGRAM_HOLE:
sim_program.append(ld_program + '\n')
elif i == self.PRE_LD_HOLE:
sim_program.append(pre_ld_program)
else:
sim_program.append(example_sim_program[i])
example_sim_program_obj.close()
sim_program_file = 't266.ptx'
sim_program_obj = open(sim_program_file, 'w')
for sim_program_line in sim_program:
sim_program_obj.write(sim_program_line)
sim_program_obj.close()
(status, output) = commands.getstatusoutput('./dryrun.out')
print status
print output
(status, output) = commands.getstatusoutput('sbatch parallel.cmd')
print status
print output
output_word = output.split()
taskTag = output_word[3]
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' +
taskTag + '.out')
while (status == 256):
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' +
taskTag + '.out')
poutput = int(output)
if (poutput < 0):
poutput = -poutput
poutput = (1 <<
(instruction_format.REG_BITS - 1)) - poutput + (
1 << (instruction_format.REG_BITS - 1))
print 'poutput: ' + str(poutput)
state[dest_text] = poutput
if self.ldAddr not in addr:
state[dest_text] = default
return state
if opcode == self.OPCODE_ST:
dmem = state['dmem']
dest_text = reg_book[dst]
st_value = state[dest_text]
self.stAddr = (self.stIMM_MASK
& instruction) >> instruction_format.IMM_BIT_BOT
self.stAddr = self.stAddr << 2
print 'store_addr' + str(self.stAddr)
outMem = ila.MemValues(instruction_format.MEM_ADDRESS_BITS,
instruction_format.DMEM_BITS, dmem.default)
for (a, v) in dmem.values:
outMem[a] = v
outMem[self.stAddr] = st_value
state['dmem'] = outMem
return state
if (opcode == self.OPCODE_MOV):
dst_text = reg_book[dst]
if (src0 >= len(reg_book)) | (dst >= len(reg_book)):
return state
if dst_text not in general_reg_book:
return state
src0_text = reg_book[src0]
if src0_text not in general_reg_book:
return state
if base:
return state
src0_data = state[src0_text]
test_program.append(op_text + '.s32 ' + dst_text + ',' +
src0_text + ';')
test_program.append('mov.s32 %r9, ' + dst_text + ';')
single_op_program = ''
for t in test_program:
single_op_program += t
single_op_program += '\n'
example_sim_program_file = 't266.ptx'
example_sim_program_obj = open(example_sim_program_file, 'r')
example_sim_program = example_sim_program_obj.readlines()
sim_program = []
for i in range(len(example_sim_program)):
if i == self.EXAMPLE_PROGRAM_HOLE:
sim_program.append(single_op_program)
else:
sim_program.append(example_sim_program[i])
example_sim_program_obj.close()
sim_program_obj = open(example_sim_program_file, 'w')
for sim_line in sim_program:
sim_program_obj.write(sim_line)
sim_program_obj.close()
(status, output) = commands.getstatusoutput('./dryrun.out')
print status
print output
(status, output) = commands.getstatusoutput('sbatch parallel.cmd')
print status
print output
output_word = output.split()
taskTag = output_word[3]
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' +
taskTag + '.out')
while (status == 256):
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' +
taskTag + '.out')
poutput = int(output)
if (poutput < 0):
poutput = -poutput
poutput = (1 <<
(instruction_format.REG_BITS - 1)) - poutput + (
1 << (instruction_format.REG_BITS - 1))
nxt_state = poutput
state[dst_text] = nxt_state
return state
dst_text = reg_book[dst]
if (src0 >= len(reg_book)) | (src1 >= len(reg_book)) | (dst >=
len(reg_book)):
return state
if dst_text not in reg_book:
return state
src0_text = reg_book[src0]
src1_text = reg_book[src1]
if (src0_text not in reg_book) | (src1_text not in reg_book):
return state
if (base):
return state
src0_data = state[src0_text]
src1_data = state[src1_text]
test_program.append(op_text + '.s32 ' + dst_text + ', ' + src0_text +
', ' + src1_text + ';')
print(op_text + ' ' + dst_text + ', ' + src0_text + ', ' + src1_text)
if dst_text in general_reg_book:
test_program.append('mov.s32 %r9, ' + dst_text + ';')
print test_program[-1]
example_sim_program_file = 't266.ptx'
example_sim_program_obj = open(example_sim_program_file, 'r')
example_sim_program = example_sim_program_obj.readlines()
sim_program = []
sim_program_first_part = 0
for i in range(len(example_sim_program)):
if i != self.EXAMPLE_PROGRAM_HOLE:
sim_program.append(example_sim_program[i])
elif i == (self.EXAMPLE_PROGRAM_HOLE + 1):
if dst_text not in general_reg_book:
sim_program.append('st.local.u32 [%rd8], %rd0;\n')
else:
sim_program.append('st.local.u32 [%rd8], %r9;\n')
else:
sim_program += test_program
sim_program += '\n'
example_sim_program_obj.close()
sim_program_obj = open(example_sim_program_file, 'w')
for sim_line in sim_program:
sim_program_obj.write(sim_line)
sim_program_obj.close()
(status, output) = commands.getstatusoutput('./dryrun.out')
print status
print output
(status, output) = commands.getstatusoutput('sbatch parallel.cmd')
print status
print output
output_word = output.split()
taskTag = output_word[3]
time.sleep(5)
(status,
output) = commands.getstatusoutput('cat slurm-' + taskTag + '.out')
while (status == 256):
time.sleep(5)
(status, output) = commands.getstatusoutput('cat slurm-' +
taskTag + '.out')
poutput = int(output)
if (poutput < 0):
poutput = -poutput
poutput = (1 << (instruction_format.REG_BITS - 1)) - poutput + (
1 << (instruction_format.REG_BITS - 1))
nxt_state = poutput
(status, output) = commands.getstatusoutput('rm a_dlin*')
(status, output) = commands.getstatusoutput('rm ' + 'slurm-' +
taskTag + '.out')
state[dst_text] = nxt_state
return state