diff options
| -rw-r--r-- | src/cuda-sim/Makefile | 156 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 140 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.h | 1 | ||||
| -rw-r--r-- | src/cuda-sim/decuda_pred_table/decuda_pred_table.cc | 79 | ||||
| -rw-r--r-- | src/cuda-sim/decuda_pred_table/decuda_pred_table.h | 1 | ||||
| -rw-r--r-- | src/cuda-sim/instructions.cc | 1185 | ||||
| -rw-r--r-- | src/cuda-sim/opcodes.def | 8 | ||||
| -rw-r--r-- | src/cuda-sim/opcodes.h | 1 | ||||
| -rw-r--r-- | src/cuda-sim/ptx.l | 24 | ||||
| -rw-r--r-- | src/cuda-sim/ptx.y | 120 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_ir.cc | 170 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_ir.h | 154 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_parser.cc | 147 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_parser.h | 14 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_sim.cc | 86 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_sim.h | 41 | ||||
| -rw-r--r-- | src/debug.cc | 2 | ||||
| -rw-r--r-- | src/gpgpu-sim/scoreboard.cc | 4 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.cc | 7 |
19 files changed, 1871 insertions, 469 deletions
diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 597f7c4..39a6641 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -91,7 +91,7 @@ ifeq ($(GNUC_CPP0X),1) endif endif -OBJS := ptx_parser.o ptx_loader.o cuda_device_printf.o instructions.o cuda-sim.o ptx_ir.o ptx_sim.o memory.o ptx-stats.o ptx.tab.o lex.ptx_.o ptxinfo.tab.o lex.ptxinfo_.o +OBJS := ptx_parser.o ptx_loader.o cuda_device_printf.o instructions.o cuda-sim.o ptx_ir.o ptx_sim.o memory.o ptx-stats.o decuda_pred_table/decuda_pred_table.o ptx.tab.o lex.ptx_.o ptxinfo.tab.o lex.ptxinfo_.o CUDART_VERSION:=$(shell nvcc --version | awk '/release/ {print $$5;}' | sed 's/,//' | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($$1), 10*$$2);}') @@ -128,7 +128,8 @@ clean: rm -f *~ *.o *.gcda *.gcno *.gcov libgpgpu_ptx_sim.a \ ptx.tab.h ptx.tab.c ptx.output lex.ptx_.c \ ptxinfo.tab.h ptxinfo.tab.c ptxinfo.output lex.ptxinfo_.c \ - instructions.h ptx_parser_decode.def directed_tests.log + instructions.h ptx_parser_decode.def directed_tests.log \ + decuda_pred_table/decuda_pred_table.o .c.o: $(CC) -c $(OPT) $< -o $*.o @@ -171,154 +172,3 @@ ptx_parser.o: ptx.tab.c ptx_parser_decode.def # DO NOT DELETE -cuda_device_printf.o: ptx_ir.h ../abstract_hardware_model.h -cuda_device_printf.o: /usr/include/assert.h /usr/include/features.h -cuda_device_printf.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h -cuda_device_printf.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h -cuda_device_printf.o: ptx_sim.h /usr/include/stdlib.h -cuda_device_printf.o: /usr/include/sys/types.h /usr/include/bits/types.h -cuda_device_printf.o: /usr/include/bits/typesizes.h /usr/include/time.h -cuda_device_printf.o: /usr/include/endian.h /usr/include/bits/endian.h -cuda_device_printf.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h -cuda_device_printf.o: /usr/include/bits/select.h /usr/include/bits/sigset.h -cuda_device_printf.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h -cuda_device_printf.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h -cuda_device_printf.o: dram_callback.h opcodes.h opcodes.def memory.h -cuda_device_printf.o: /usr/include/string.h /usr/include/stdio.h -cuda_device_printf.o: /usr/include/libio.h /usr/include/_G_config.h -cuda_device_printf.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h -cuda_device_printf.o: /usr/include/bits/sys_errlist.h ../option_parser.h -cuda-sim.o: cuda-sim.h ../abstract_hardware_model.h /usr/include/stdlib.h -cuda-sim.o: /usr/include/features.h /usr/include/sys/cdefs.h -cuda-sim.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h -cuda-sim.o: /usr/include/gnu/stubs-64.h /usr/include/sys/types.h -cuda-sim.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h -cuda-sim.o: /usr/include/time.h /usr/include/endian.h -cuda-sim.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h -cuda-sim.o: /usr/include/sys/select.h /usr/include/bits/select.h -cuda-sim.o: /usr/include/bits/sigset.h /usr/include/bits/time.h -cuda-sim.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h -cuda-sim.o: /usr/include/alloca.h ptx_ir.h /usr/include/assert.h ptx_sim.h -cuda-sim.o: dram_callback.h opcodes.h opcodes.def memory.h -cuda-sim.o: /usr/include/string.h /usr/include/stdio.h /usr/include/libio.h -cuda-sim.o: /usr/include/_G_config.h /usr/include/wchar.h -cuda-sim.o: /usr/include/bits/stdio_lim.h /usr/include/bits/sys_errlist.h -cuda-sim.o: ../option_parser.h ../intersim/statwraper.h ptx-stats.h -cuda-sim.o: ptx_loader.h ptx_parser.h -instructions.o: ptx_ir.h ../abstract_hardware_model.h /usr/include/assert.h -instructions.o: /usr/include/features.h /usr/include/sys/cdefs.h -instructions.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h -instructions.o: /usr/include/gnu/stubs-64.h ptx_sim.h /usr/include/stdlib.h -instructions.o: /usr/include/sys/types.h /usr/include/bits/types.h -instructions.o: /usr/include/bits/typesizes.h /usr/include/time.h -instructions.o: /usr/include/endian.h /usr/include/bits/endian.h -instructions.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h -instructions.o: /usr/include/bits/select.h /usr/include/bits/sigset.h -instructions.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h -instructions.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h -instructions.o: dram_callback.h opcodes.h opcodes.def memory.h -instructions.o: /usr/include/string.h /usr/include/stdio.h -instructions.o: /usr/include/libio.h /usr/include/_G_config.h -instructions.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h -instructions.o: /usr/include/bits/sys_errlist.h ../option_parser.h -instructions.o: /usr/include/math.h /usr/include/bits/huge_val.h -instructions.o: /usr/include/bits/mathdef.h /usr/include/bits/mathcalls.h -instructions.o: /usr/include/fenv.h /usr/include/bits/fenv.h cuda-math.h -instructions.o: ptx_loader.h -memory.o: memory.h ../abstract_hardware_model.h /usr/include/assert.h -memory.o: /usr/include/features.h /usr/include/sys/cdefs.h -memory.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h -memory.o: /usr/include/gnu/stubs-64.h /usr/include/string.h -memory.o: /usr/include/stdio.h /usr/include/bits/types.h -memory.o: /usr/include/bits/typesizes.h /usr/include/libio.h -memory.o: /usr/include/_G_config.h /usr/include/wchar.h -memory.o: /usr/include/bits/stdio_lim.h /usr/include/bits/sys_errlist.h -memory.o: /usr/include/stdlib.h /usr/include/sys/types.h /usr/include/time.h -memory.o: /usr/include/endian.h /usr/include/bits/endian.h -memory.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h -memory.o: /usr/include/bits/select.h /usr/include/bits/sigset.h -memory.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h -memory.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h ../debug.h -memory.o: ../abstract_hardware_model.h -ptx_ir.o: ptx_parser.h ../abstract_hardware_model.h ptx_ir.h -ptx_ir.o: /usr/include/assert.h /usr/include/features.h -ptx_ir.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h -ptx_ir.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h ptx_sim.h -ptx_ir.o: /usr/include/stdlib.h /usr/include/sys/types.h -ptx_ir.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h -ptx_ir.o: /usr/include/time.h /usr/include/endian.h -ptx_ir.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h -ptx_ir.o: /usr/include/sys/select.h /usr/include/bits/select.h -ptx_ir.o: /usr/include/bits/sigset.h /usr/include/bits/time.h -ptx_ir.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h -ptx_ir.o: /usr/include/alloca.h dram_callback.h opcodes.h opcodes.def -ptx_ir.o: memory.h /usr/include/string.h /usr/include/stdio.h -ptx_ir.o: /usr/include/libio.h /usr/include/_G_config.h /usr/include/wchar.h -ptx_ir.o: /usr/include/bits/stdio_lim.h /usr/include/bits/sys_errlist.h -ptx_ir.o: ../option_parser.h cuda-sim.h -ptx_loader.o: ptx_loader.h ptx_ir.h ../abstract_hardware_model.h -ptx_loader.o: /usr/include/assert.h /usr/include/features.h -ptx_loader.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h -ptx_loader.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h ptx_sim.h -ptx_loader.o: /usr/include/stdlib.h /usr/include/sys/types.h -ptx_loader.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h -ptx_loader.o: /usr/include/time.h /usr/include/endian.h -ptx_loader.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h -ptx_loader.o: /usr/include/sys/select.h /usr/include/bits/select.h -ptx_loader.o: /usr/include/bits/sigset.h /usr/include/bits/time.h -ptx_loader.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h -ptx_loader.o: /usr/include/alloca.h dram_callback.h opcodes.h opcodes.def -ptx_loader.o: memory.h /usr/include/string.h /usr/include/stdio.h -ptx_loader.o: /usr/include/libio.h /usr/include/_G_config.h -ptx_loader.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h -ptx_loader.o: /usr/include/bits/sys_errlist.h ../option_parser.h cuda-sim.h -ptx_loader.o: ptx_parser.h /usr/include/dirent.h /usr/include/bits/dirent.h -ptx_loader.o: /usr/include/bits/posix1_lim.h /usr/include/bits/local_lim.h -ptx_loader.o: /usr/include/linux/limits.h -ptx_parser.o: ptx_parser.h ../abstract_hardware_model.h ptx_ir.h -ptx_parser.o: /usr/include/assert.h /usr/include/features.h -ptx_parser.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h -ptx_parser.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h ptx_sim.h -ptx_parser.o: /usr/include/stdlib.h /usr/include/sys/types.h -ptx_parser.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h -ptx_parser.o: /usr/include/time.h /usr/include/endian.h -ptx_parser.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h -ptx_parser.o: /usr/include/sys/select.h /usr/include/bits/select.h -ptx_parser.o: /usr/include/bits/sigset.h /usr/include/bits/time.h -ptx_parser.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h -ptx_parser.o: /usr/include/alloca.h dram_callback.h opcodes.h opcodes.def -ptx_parser.o: memory.h /usr/include/string.h /usr/include/stdio.h -ptx_parser.o: /usr/include/libio.h /usr/include/_G_config.h -ptx_parser.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h -ptx_parser.o: /usr/include/bits/sys_errlist.h ../option_parser.h -ptx_sim.o: ptx_sim.h /usr/include/stdlib.h /usr/include/features.h -ptx_sim.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h -ptx_sim.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h -ptx_sim.o: /usr/include/sys/types.h /usr/include/bits/types.h -ptx_sim.o: /usr/include/bits/typesizes.h /usr/include/time.h -ptx_sim.o: /usr/include/endian.h /usr/include/bits/endian.h -ptx_sim.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h -ptx_sim.o: /usr/include/bits/select.h /usr/include/bits/sigset.h -ptx_sim.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h -ptx_sim.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h -ptx_sim.o: dram_callback.h ../abstract_hardware_model.h /usr/include/assert.h -ptx_sim.o: opcodes.h opcodes.def ptx_ir.h memory.h /usr/include/string.h -ptx_sim.o: /usr/include/stdio.h /usr/include/libio.h /usr/include/_G_config.h -ptx_sim.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h -ptx_sim.o: /usr/include/bits/sys_errlist.h ../option_parser.h -ptx-stats.o: ptx_ir.h ../abstract_hardware_model.h /usr/include/assert.h -ptx-stats.o: /usr/include/features.h /usr/include/sys/cdefs.h -ptx-stats.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h -ptx-stats.o: /usr/include/gnu/stubs-64.h ptx_sim.h /usr/include/stdlib.h -ptx-stats.o: /usr/include/sys/types.h /usr/include/bits/types.h -ptx-stats.o: /usr/include/bits/typesizes.h /usr/include/time.h -ptx-stats.o: /usr/include/endian.h /usr/include/bits/endian.h -ptx-stats.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h -ptx-stats.o: /usr/include/bits/select.h /usr/include/bits/sigset.h -ptx-stats.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h -ptx-stats.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h -ptx-stats.o: dram_callback.h opcodes.h opcodes.def memory.h -ptx-stats.o: /usr/include/string.h /usr/include/stdio.h /usr/include/libio.h -ptx-stats.o: /usr/include/_G_config.h /usr/include/wchar.h -ptx-stats.o: /usr/include/bits/stdio_lim.h /usr/include/bits/sys_errlist.h -ptx-stats.o: ../option_parser.h ptx-stats.h diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 2c317b2..0a6ccd6 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -82,6 +82,7 @@ #include "../gpgpu-sim/gpu-sim.h" #include "ptx_sim.h" #include "../gpgpusim_entrypoint.h" +#include "decuda_pred_table/decuda_pred_table.h" int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; @@ -90,6 +91,9 @@ int g_ptx_kernel_count = -1; // used for classification stat collection purposes int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; +// Output debug information to file options +FILE* ptx_inst_debug_file; + unsigned g_ptx_sim_num_insn = 0; std::map<const struct textureReference*,const struct cudaArray*> TextureToArrayMap; // texture bindings std::map<const struct textureReference*, const struct textureInfo*> TextureToInfoMap; @@ -245,7 +249,7 @@ void function_info::ptx_assemble() g_assemble_code_next_pc=PC; for ( unsigned ii=0; ii < n; ii++ ) { // handle branch instructions ptx_instruction *pI = m_instr_mem[ii]; - if ( pI->get_opcode() == BRA_OP ) { + if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP ) { operand_info &target = pI->dst(); //get operand, e.g. target name if ( labels.find(target.name()) == labels.end() ) { printf("GPGPU-Sim PTX: Loader error (%s:%u): Branch label \"%s\" does not appear in assembly code.", @@ -263,17 +267,28 @@ void function_info::ptx_assemble() create_basic_blocks(); connect_basic_blocks(); + bool modified = false; + do { + find_dominators(); + find_idominators(); + modified = connect_break_targets(); + } while (modified == true); + if ( g_debug_execution>=50 ) { print_basic_blocks(); print_basic_block_links(); print_basic_block_dot(); } + if ( g_debug_execution>=2 ) { + print_dominators(); + } find_postdominators(); find_ipostdominators(); if ( g_debug_execution>=50 ) { print_postdominators(); print_ipostdominators(); } + m_assembled = true; } @@ -534,6 +549,8 @@ static void get_opcode_info( const ptx_instruction *pI, unsigned opcode, unsigne *op_type = STORE_OP; } else if ( opcode == BRA_OP ) { *op_type = BRANCH_OP; + } else if ( opcode == BREAKADDR_OP ) { + *op_type = BRANCH_OP; } else if ( opcode == TEX_OP ) { *op_type = LOAD_OP; } else if ( opcode == ATOM_OP ) { @@ -596,6 +613,17 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread, unsigned cycles; get_opcode_info(pI,opcode,&cycles,op_type); + // Quick fix for memory operands in ALU instructions + if( pI->has_memory_read() ) + *op_type = LOAD_OP; + else if( pI->has_memory_write() ) + *op_type = STORE_OP; + + if( pI->has_memory_read() && pI->has_memory_write() ) { + printf("Instruction has both a memory read and memory write - not supported by timing simulator."); + assert(0); + } + // Get register operands int n=0,m=0; ptx_instruction::const_iterator op=pI->op_iter_begin(); @@ -659,6 +687,35 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread, if(o.is_memory_operand2()) { // memory operand with one address register (ex. g[$r1] or s[$r2+0x4]) if(o.get_double_operand_type() == 0 && o.is_memory_operand()){ + *ar1 = o.reg_num(); + } + // memory operand with two address register (ex. s[$r1+$r1] or g[$r1+=$r2]) + else if(o.get_double_operand_type() == 1 || o.get_double_operand_type() == 2) { + *ar1 = o.reg1_num(); + *ar2 = o.reg2_num(); + } + } + } + } + + // Get predicate + if(pI->has_pred()) { + const operand_info &p = pI->get_pred(); + *pred = p.reg_num(); + } + + // Get address registers inside memory operands. + // Assuming only one memory operand per instruction, + // and maximum of two address registers for one memory operand. + if( pI->has_memory_read() || pI->has_memory_write() ) { + ptx_instruction::const_iterator op=pI->op_iter_begin(); + for ( ; op != pI->op_iter_end(); op++, n++ ) { //process operands + const operand_info &o = *op; + + // memory operand with addressing (ex. s[0x4] or g[$r1]) + if(o.is_memory_operand2()) { + // memory operand with one address register (ex. g[$r1] or s[$r2+0x4]) + if(o.get_double_operand_type() == 0 && o.is_memory_operand()){ const symbol *base_addr = o.get_symbol(); if( base_addr->is_reg() ) *ar1 = base_addr->reg_num(); @@ -700,6 +757,7 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg tmp.offset = args->m_offset; tmp.type = 0; i->second.add_data(tmp); + i->second.add_offset((unsigned) args->m_offset); } else { // This should only happen for OpenCL: // @@ -778,6 +836,30 @@ void function_info::finalize( memory_space *param_mem ) } } +void function_info::param_to_shared( memory_space *shared_mem, symbol_table *symtab ) +{ +/*copies parameters into simulated shared memory*/ + for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { + param_info &p = i->second; + std::string name = p.get_name(); + int type = p.get_type(); + param_t value = p.get_value(); + value.type = type; + symbol *param = symtab->lookup(name.c_str()); + unsigned xtype = param->type()->get_key().scalar_type(); + assert(xtype==(unsigned)type); + + int tmp; + size_t size; + unsigned offset = p.get_offset(); + type_info_key::type_decode(xtype,size,tmp); + + // Write to shared memory - offset + 0x10 + shared_mem->write(offset+0x10,size/8,&value,NULL,NULL); + } +} + + void function_info::list_param( FILE *fout ) const { for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -827,17 +909,19 @@ unsigned datatype2size( unsigned data_type ) case F32_TYPE: data_size = 4; break; case B64_TYPE: + case BB64_TYPE: case S64_TYPE: case U64_TYPE: case F64_TYPE: data_size = 8; break; + case BB128_TYPE: + data_size = 16; break; default: assert(0); break; } return data_size; } unsigned g_warp_active_mask; -FILE *ptx_inst_debug_file; void function_info::ptx_exec_inst( ptx_thread_info *thread, addr_t *addr, @@ -858,6 +942,11 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, thread->clearRPC(); thread->m_last_set_operand_value.u64 = 0; + if(thread->is_done()) + { + printf("attempted to execute instruction on a thread that is already done.\n"); + assert(0); + } if ( g_debug_execution >= 6 ) { if ( (g_debug_thread_uid==0) || (thread->get_uid() == (unsigned)g_debug_thread_uid) ) { thread->clear_modifiedregs(); @@ -866,8 +955,12 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, } if( pI->has_pred() ) { const operand_info &pred = pI->get_pred(); - ptx_reg_t pred_value = thread->get_operand_value(pred); - skip = !pred_value.pred ^ pI->get_pred_neg(); + ptx_reg_t pred_value = thread->get_operand_value(pred, pred, PRED_TYPE, thread, 0); + if(pI->get_pred_mod() == -1) { + skip = (pred_value.pred & 0x0001) ^ pI->get_pred_neg(); //ptxplus inverts the zero flag + } else { + skip = !pred_lookup(pI->get_pred_mod(), pred_value.pred & 0x000F); + } } g_warp_active_mask = warp_active_mask; if( !skip ) { @@ -875,10 +968,28 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,thread); op_classification = CLASSIFICATION; break; #include "opcodes.def" #undef OP_DEF - default: - printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); - break; + + default: + printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); + break; } + + // Run exit instruction if exit option included + if(pI->is_exit()) + exit_impl(pI,thread); + } + + // Output instruction information to file and stdout + if( g_ptx_inst_debug_to_file != 0 && + (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == thread->get_uid()) ) { + dim3 ctaid = thread->get_ctaid(); + dim3 tid = thread->get_tid(); + fprintf(ptx_inst_debug_file, + "[thd=%u] : (%s:%u - %s)\n", + thread->get_uid(), + pI->source_file(), pI->source_line(), pI->get_source() ); + //fprintf(ptx_inst_debug_file, "has memory read=%d, has memory write=%d\n", pI->has_memory_read(), pI->has_memory_write()); + fflush(ptx_inst_debug_file); } if ( ptx_debug_exec_dump_cond<5>(thread->get_uid(), pc) ) { @@ -958,15 +1069,15 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread, if( g_ptx_inst_debug_to_file != 0 && (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == thread->get_uid()) ) { thread->dump_modifiedregs(ptx_inst_debug_file); - thread->dump_regs(ptx_inst_debug_file); } if ( g_debug_execution >= 6 ) { if ( ptx_debug_exec_dump_cond<6>(thread->get_uid(), pc) ) - thread->dump_modifiedregs(); - } else if ( g_debug_execution >= 10 ) { + thread->dump_modifiedregs(stdout); + } + if ( g_debug_execution >= 10 ) { if ( ptx_debug_exec_dump_cond<10>(thread->get_uid(), pc) ) - thread->dump_regs(); + thread->dump_regs(stdout); } thread->update_pc(); g_ptx_sim_num_insn++; @@ -1168,6 +1279,8 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid thd->set_ntid(g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z); thd->set_ctaid(g_gx,g_gy,g_gz); thd->set_tid(tx,ty,tz); + if( g_entrypoint_func_info->get_ptx_version().extensions() ) + thd->cpy_tid_to_reg(tx,ty,tz); thd->set_hw_tid((unsigned)-1); thd->set_hw_wid((unsigned)-1); thd->set_hw_ctaid((unsigned)-1); @@ -1175,6 +1288,9 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid thd->set_hw_sid((unsigned)-1); thd->set_valid(); thd->m_shared_mem = shared_mem; + function_info *finfo = thd->func_info(); + symbol_table *st = finfo->get_symtab(); + thd->func_info()->param_to_shared(thd->m_shared_mem,st); thd->m_cta_info = cta_info; cta_info->add_thread(thd); thd->m_local_mem = local_mem; @@ -1691,7 +1807,7 @@ void ptx_dump_regs( void *thd ) if ( thd == NULL ) return; ptx_thread_info *t = (ptx_thread_info *) thd; - t->dump_regs(); + t->dump_regs(stdout); } unsigned ptx_set_tex_cache_linesize(unsigned linesize) diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 6f931fc..d7d3f8a 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -20,6 +20,7 @@ extern std::map<std::string,function_info*> *g_kernel_name_to_function_lookup; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; extern int g_ptx_kernel_count; // used for classification stat collection purposes +extern FILE* ptx_inst_debug_file; extern void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, diff --git a/src/cuda-sim/decuda_pred_table/decuda_pred_table.cc b/src/cuda-sim/decuda_pred_table/decuda_pred_table.cc new file mode 100644 index 0000000..5435d00 --- /dev/null +++ b/src/cuda-sim/decuda_pred_table/decuda_pred_table.cc @@ -0,0 +1,79 @@ +/*Copyright (c) 2007, Wladimir J. van der Laan + +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + + * Redistributions of source code must retain the above copyright notice, + this list of conditions and the following disclaimer. + * Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimer in the + documentation and/or other materials provided with the distribution. + * Neither the name of the organization nor the names of its + contributors may be used to endorse or promote products derived from this + software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS +"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT +LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR +A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR +CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF +LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING +NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.*/ + +#include "decuda_pred_table.h" + +bool pred_lookup(int condition, int flags) +{ + + // Logic table for G80 architecture, all condition codes against all + // flag combinations. This was evaluated on actual hardware. + // The flags are assigned to values like this: + // 1 Z zero flag + // 2 S sign flag + // 4 C carry flag + // 8 O overflow flag + + // + // fl 0 1 2 3 4 5 6 7 8 9 a b c d e f + static bool const pred_table[32][16] = + {{0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, // 00 fl + {0, 0, 1, 0, 0, 0, 1, 0, 1, 1, 0, 1, 1, 1, 0, 1}, // 01 lt + {0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0}, // 02 eq + {0, 1, 1, 0, 0, 1, 1, 0, 1, 1, 0, 0, 1, 1, 0, 0}, // 03 le + {1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0}, // 04 gt + {1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0}, // 05 ne (also nz) + {1, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 0, 1, 1}, // 06 ge + {1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1, 0}, // 07 leg + {0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1}, // 08 nan + {0, 0, 1, 1, 0, 0, 1, 1, 1, 1, 0, 0, 1, 1, 0, 0}, // 09 ltu + {0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1}, // 0a equ (also zf) + {0, 1, 1, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0, 1}, // 0b leu + {1, 0, 0, 1, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 1, 1}, // 0c gtu + {1, 0, 1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1, 0, 1, 1}, // 0d neu + {1, 1, 0, 1, 1, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0}, // 0e geu + {1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}, // 0f tr + {0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1}, // 10 of + {0, 0, 0, 0, 1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1}, // 11 cf + {0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 1, 0, 1, 0}, // 12 ab + {0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1}, // 13 sf + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, + {1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0, 1, 1, 0, 0}, // 1c nsf + {1, 1, 1, 1, 0, 1, 0, 1, 1, 1, 1, 1, 0, 1, 0, 1}, // 1d ble + {1, 1, 1, 1, 0, 0, 0, 0, 1, 1, 1, 1, 0, 0, 0, 0}, // 1e ncf + {1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0}}; // 1f nof + + return pred_table[condition][flags]; +} diff --git a/src/cuda-sim/decuda_pred_table/decuda_pred_table.h b/src/cuda-sim/decuda_pred_table/decuda_pred_table.h new file mode 100644 index 0000000..c2b827c --- /dev/null +++ b/src/cuda-sim/decuda_pred_table/decuda_pred_table.h @@ -0,0 +1 @@ +bool pred_lookup(int condition, int flags); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index cbf4c44..d9a9630 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -95,11 +95,27 @@ extern std::map<struct textureReference*,struct textureInfo*> TextureToInfoMap; extern std::map<std::string, struct textureReference*> NameToTextureMap; void inst_not_implemented( const ptx_instruction * pI ) ; +ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread); unsigned unfound_register_warned = 0; -ptx_reg_t ptx_thread_info::get_operand_value( const symbol *reg ) +void sign_extend( ptx_reg_t &data, unsigned src_size, const operand_info &dst ); + +void ptx_thread_info::set_reg( const symbol *reg, const ptx_reg_t &value ) +{ + assert( reg != NULL ); + if( reg->name() == "_" ) return; + assert( !m_regs.empty() ); + assert( reg->uid() > 0 ); + m_regs.back()[ reg ] = value; + if (m_enable_debug_trace ) + m_debug_trace_regs_modified.back()[ reg ] = value; + m_last_set_operand_value = value; +} + +ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) { - // assume that the given symbol is a register and try to find it in the register hash map + assert( reg != NULL ); + assert( !m_regs.empty() ); reg_map_t::iterator regs_iter = m_regs.back().find(reg); if (regs_iter == m_regs.back().end()) { assert( reg->type()->get_key().is_reg() ); @@ -107,7 +123,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const symbol *reg ) unsigned call_uid = m_callstack.back().m_call_uid; ptx_reg_t uninit_reg; uninit_reg.u32 = 0xDEADBEEF; - set_operand_value(reg, uninit_reg); // give it a value since we are going to warn the user anyway + set_reg(reg, uninit_reg); // give it a value since we are going to warn the user anyway std::string file_loc = get_location(); if( !unfound_register_warned ) { printf("GPGPU-Sim PTX: WARNING (%s) ** reading undefined register \'%s\' (cuid:%u). Setting to 0XDEADBEEF.\n", @@ -116,67 +132,203 @@ ptx_reg_t ptx_thread_info::get_operand_value( const symbol *reg ) } regs_iter = m_regs.back().find(reg); } + if (m_enable_debug_trace ) + m_debug_trace_regs_read.back()[ reg ] = regs_iter->second; return regs_iter->second; } -ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op ) +ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_info dstInfo, unsigned opType, ptx_thread_info *thread, int derefFlag ) { - ptx_reg_t result; - const char *name = NULL; - if ( op.is_reg() ) { - result = get_operand_value( op.get_symbol() ); - } else if ( op.is_builtin()) { - result = get_builtin( op.get_int(), op.get_addr_offset() ); - } else if ( op.is_memory_operand() ) { - // a few options here... - const symbol *sym = op.get_symbol(); - const type_info *type = sym->type(); - const type_info_key &info = type->get_key(); + ptx_reg_t result, tmp; + + if(op.get_double_operand_type() == 0) { + if(((opType != BB128_TYPE) && (opType != BB64_TYPE)) || (op.get_addr_space() != 0)) { + if ( op.is_reg() ) { + result = get_reg( op.get_symbol() ); + } else if ( op.is_builtin()) { + result = get_builtin( op.get_int(), op.get_addr_offset() ); + } else if ( op.is_memory_operand() ) { + // a few options here... + const symbol *sym = op.get_symbol(); + const type_info *type = sym->type(); + const type_info_key &info = type->get_key(); + + if ( info.is_reg() ) { + const symbol *name = op.get_symbol(); + result.u64 = get_reg(name).u64 + op.get_addr_offset(); + } else if ( info.is_param_kernel() ) { + result = sym->get_address() + op.get_addr_offset(); + } else if ( info.is_param_local() ) { + result = sym->get_address() + op.get_addr_offset(); + } else if ( info.is_global() ) { + assert( op.get_addr_offset() == 0 ); + result = sym->get_address(); + } else if ( info.is_local() ) { + result = sym->get_address() + op.get_addr_offset(); + } else if ( info.is_const() ) { + result = sym->get_address() + op.get_addr_offset(); + } else if ( op.is_shared() ) { + result = op.get_symbol()->get_address() + op.get_addr_offset(); + } else { + const char *name = op.name().c_str(); + printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown memory operand type for %s\n", name ); + abort(); + } - if ( info.is_reg() ) { - name = op.name().c_str(); - reg_map_t::iterator regs_iter = m_regs.back().find(sym); - assert( regs_iter != m_regs.back().end() ); - ptx_reg_t baseaddr = regs_iter->second; - result.u64 = baseaddr.u64 + op.get_addr_offset(); - } else if ( info.is_param_kernel() ) { - result = sym->get_address() + op.get_addr_offset(); - } else if ( info.is_param_local() ) { - result = sym->get_address() + op.get_addr_offset(); - } else if ( info.is_global() ) { - assert( op.get_addr_offset() == 0 ); - result = sym->get_address(); - } else if ( info.is_local() ) { - result = sym->get_address() + op.get_addr_offset(); - } else if ( info.is_const() ) { - result = sym->get_address() + op.get_addr_offset(); - } else if ( op.is_shared() ) { - result = op.get_symbol()->get_address() + op.get_addr_offset(); + } else if ( op.is_literal() ) { + result = op.get_literal_value(); + } else if ( op.is_label() ) { + result = op.get_symbol()->get_address(); + } else if ( op.is_shared() ) { + result = op.get_symbol()->get_address(); + } else if ( op.is_const() ) { + result = op.get_symbol()->get_address(); + } else if ( op.is_global() ) { + result = op.get_symbol()->get_address(); + } else if ( op.is_local() ) { + result = op.get_symbol()->get_address(); + } else { + const char *name = op.name().c_str(); + printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); + assert(0); + } + + if(op.get_operand_lohi() == 1) + result.u64 = result.u64 & 0xFFFF; + else if(op.get_operand_lohi() == 2) + result.u64 = (result.u64>>16) & 0xFFFF; + } else if (opType == BB128_TYPE) { + // b128 + result.u128.lowest = get_reg( op.vec_symbol(0) ).u32; + result.u128.low = get_reg( op.vec_symbol(1) ).u32; + result.u128.high = get_reg( op.vec_symbol(2) ).u32; + result.u128.highest = get_reg( op.vec_symbol(3) ).u32; } else { - name = op.name().c_str(); - printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown memory operand type for %s\n", name ); - abort(); + // bb64 + result.bits.ls = get_reg( op.vec_symbol(0) ).u32; + result.bits.ms = get_reg( op.vec_symbol(1) ).u32; } + } else if (op.get_double_operand_type() == 1) { + ptx_reg_t firstHalf, secondHalf; + firstHalf = get_reg( op.vec_symbol(0) ).u64; + secondHalf = get_reg( op.vec_symbol(1) ).u64; + if(op.get_operand_lohi() == 1) + secondHalf.u64 = secondHalf.u64 & 0xFFFF; + else if(op.get_operand_lohi() == 2) + secondHalf.u64 = (secondHalf.u64>>16) & 0xFFFF; + result = firstHalf.u64 + secondHalf.u64; + } else if (op.get_double_operand_type() == 2) { + // s[reg1 += reg2] + // reg1 is incremented after value is returned: the value returned is s[reg1] + ptx_reg_t firstHalf, secondHalf; + firstHalf = get_reg(op.vec_symbol(0)).u64; + secondHalf = get_reg(op.vec_symbol(1)).u64; + if(op.get_operand_lohi() == 1) + secondHalf.u64 = secondHalf.u64 & 0xFFFF; + else if(op.get_operand_lohi() == 2) + secondHalf.u64 = (secondHalf.u64>>16) & 0xFFFF; + result = firstHalf.u64; + firstHalf = firstHalf.u64 + secondHalf.u64; + set_reg(op.vec_symbol(0),firstHalf); + } else if (op.get_double_operand_type() == 3) { + // s[reg += immediate] + // reg is incremented after value is returned: the value returned is s[reg] + ptx_reg_t firstHalf; + firstHalf = get_reg(op.get_symbol()).u64; + result = firstHalf.u64; + firstHalf = firstHalf.u64 + op.get_addr_offset(); + set_reg(op.get_symbol(),firstHalf); + } + + ptx_reg_t finalResult; + memory_space *mem = NULL; + size_t size; + int t; + finalResult.u64=0; + + //complete other cases for reading from memory, such as reading from other const memory + if((op.get_addr_space() == 1)&&(derefFlag)) { + // global memory - g[4], g[$r0] + mem = g_global_mem; + type_info_key::type_decode(opType,size,t); + mem->read(result.u32,size/8,&finalResult.u128); + thread->m_last_effective_address = result.u32; + thread->m_last_memory_space = global_space; + + if( opType == S16_TYPE || opType == S32_TYPE ) + sign_extend(finalResult,size,dstInfo); + } else if((op.get_addr_space() == 2)&&(derefFlag)) { + // shared memory - s[4], s[$r0] + mem = thread->m_shared_mem; + type_info_key::type_decode(opType,size,t); + mem->read(result.u32,size/8,&finalResult.u128); + thread->m_last_effective_address = result.u32; + thread->m_last_memory_space = shared_space; - } else if ( op.is_literal() ) { - result = op.get_literal_value(); - } else if ( op.is_label() ) { - result = op.get_symbol()->get_address(); - } else if ( op.is_shared() ) { - result = op.get_symbol()->get_address(); - } else if ( op.is_const() ) { - result = op.get_symbol()->get_address(); - } else if ( op.is_global() ) { - result = op.get_symbol()->get_address(); - } else if ( op.is_local() ) { - result = op.get_symbol()->get_address(); + if( opType == S16_TYPE || opType == S32_TYPE ) + sign_extend(finalResult,size,dstInfo); + } else if((op.get_addr_space() == 3)&&(derefFlag)) { + // const memory - ce0c1[4], ce0c1[$r0] + mem = g_global_mem; + type_info_key::type_decode(opType,size,t); + mem->read((result.u32 + op.get_const_mem_offset()),size/8,&finalResult.u128); + thread->m_last_effective_address = result.u32; + thread->m_last_memory_space = const_space; + if( opType == S16_TYPE || opType == S32_TYPE ) + sign_extend(finalResult,size,dstInfo); + } else if((op.get_addr_space() == 4)&&(derefFlag)) { + // local memory - l0[4], l0[$r0] + mem = thread->m_local_mem; + type_info_key::type_decode(opType,size,t); + mem->read(result.u32,size/8,&finalResult.u128); + thread->m_last_effective_address = result.u32; + thread->m_last_memory_space = local_space; + if( opType == S16_TYPE || opType == S32_TYPE ) + sign_extend(finalResult,size,dstInfo); } else { - name = op.name().c_str(); - printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); - assert(0); + finalResult = result; } - return result; + if((op.get_operand_neg() == true)&&(derefFlag)) { + switch( opType ) { + // Default to f32 for now, need to add support for others + case S8_TYPE: + case U8_TYPE: + case B8_TYPE: + finalResult.s8 = -finalResult.s8; + break; + case S16_TYPE: + case U16_TYPE: + case B16_TYPE: + finalResult.s16 = -finalResult.s16; + break; + case S32_TYPE: + case U32_TYPE: + case B32_TYPE: + finalResult.s32 = -finalResult.s32; + break; + case S64_TYPE: + case U64_TYPE: + case B64_TYPE: + finalResult.s64 = -finalResult.s64; + break; + case F16_TYPE: + finalResult.f16 = -finalResult.f16; + break; + case F32_TYPE: + finalResult.f32 = -finalResult.f32; + break; + case F64_TYPE: + finalResult.f64 = -finalResult.f64; + break; + default: + assert(0); + } + + } + + return finalResult; + } unsigned get_operand_nbits( const operand_info &op ) @@ -218,9 +370,11 @@ void ptx_thread_info::get_vector_operand_values( const operand_info &op, ptx_reg for (int idx = num_elements - 1; idx >= 0; --idx) { const symbol *sym = NULL; sym = op.vec_symbol(idx); - reg_map_t::iterator reg_iter = m_regs.back().find(sym); - assert( reg_iter != m_regs.back().end() ); - ptx_regs[idx] = reg_iter->second; + if( strcmp(sym->name().c_str(),"_") != 0) { + reg_map_t::iterator reg_iter = m_regs.back().find(sym); + assert( reg_iter != m_regs.back().end() ); + ptx_regs[idx] = reg_iter->second; + } } } @@ -246,22 +400,251 @@ void sign_extend( ptx_reg_t &data, unsigned src_size, const operand_info &dst ) data.u64 |= mask; } -void ptx_thread_info::set_operand_value( const operand_info &dst, const ptx_reg_t &data ) +void ptx_thread_info::set_operand_value( const operand_info &dst, const ptx_reg_t &data, unsigned type, ptx_thread_info *thread, const ptx_instruction *pI, int overflow, int carry ) { - m_regs.back()[ dst.get_symbol() ] = data; - if (m_enable_debug_trace ) { - m_debug_trace_regs_modified[ dst.get_symbol() ] = data; - } - m_last_set_operand_value = data; + thread->set_operand_value( dst, data, type, thread, pI ); + + if (dst.get_double_operand_type() == -2) + { + ptx_reg_t predValue; + + const symbol *sym = dst.vec_symbol(0); + predValue.u64 = (m_regs.back()[ sym ].u64) & ~(0x0C); + predValue.u64 |= ((overflow & 0x01)<<3); + predValue.u64 |= ((carry & 0x01)<<2); + + set_reg(sym,predValue); + } + else if (dst.get_double_operand_type() == 0) + { + //intentionally do nothing + } + else + { + printf("Unexpected double destination\n"); + assert(0); + } + } -void ptx_thread_info::set_operand_value( const symbol *dst, const ptx_reg_t &data ) +void ptx_thread_info::set_operand_value( const operand_info &dst, const ptx_reg_t &data, unsigned type, ptx_thread_info *thread, const ptx_instruction *pI ) { - m_regs.back()[ dst ] = data; - if (m_enable_debug_trace ) { - m_debug_trace_regs_modified[ dst ] = data; + ptx_reg_t dstData; + memory_space *mem = NULL; + size_t size; + int t; + + type_info_key::type_decode(type,size,t); + + /*complete this section for other cases*/ + if(dst.get_addr_space() == 0) + { + ptx_reg_t setValue; + setValue.u64 = data.u64; + + // Double destination in set instruction ($p0|$p1) - second is negation of first + if (dst.get_double_operand_type() == -1) + { + ptx_reg_t setValue2; + const symbol *name1 = dst.vec_symbol(0); + const symbol *name2 = dst.vec_symbol(1); + + if ( (type==F16_TYPE)||(type==F32_TYPE)||(type==F64_TYPE) ) { + setValue2.f32 = (setValue.u64==0)?1.0f:0.0f; + } else { + setValue2.u32 = (setValue.u64==0)?0xFFFFFFFF:0; + } + + set_reg(name1,setValue); + set_reg(name2,setValue2); + } + + // Double destination in cvt,shr,mul,etc. instruction ($p0|$r4) - second register operand receives data, first predicate operand + // is set as $p0=($r4!=0) + // Also for Double destination in set instruction ($p0/$r1) + else if ((dst.get_double_operand_type() == -2)||(dst.get_double_operand_type() == -3)) + { + ptx_reg_t predValue; + const symbol *predName = dst.vec_symbol(0); + const symbol *regName = dst.vec_symbol(1); + predValue.u64 = 0; + + switch ( type ) { + case S8_TYPE: + if((setValue.s8 & 0x7F) == 0) + predValue.u64 |= 1; + break; + case S16_TYPE: + if((setValue.s16 & 0x7FFF) == 0) + predValue.u64 |= 1; + break; + case S32_TYPE: + if((setValue.s32 & 0x7FFFFFFF) == 0) + predValue.u64 |= 1; + break; + case S64_TYPE: + if((setValue.s64 & 0x7FFFFFFFFFFFFFFF) == 0) + predValue.u64 |= 1; + break; + case U8_TYPE: + case B8_TYPE: + if(setValue.u8 == 0) + predValue.u64 |= 1; + break; + case U16_TYPE: + case B16_TYPE: + if(setValue.u16 == 0) + predValue.u64 |= 1; + break; + case U32_TYPE: + case B32_TYPE: + if(setValue.u32 == 0) + predValue.u64 |= 1; + break; + case U64_TYPE: + case B64_TYPE: + if(setValue.u64 == 0) + predValue.u64 |= 1; + break; + case F16_TYPE: + if(setValue.f16 == 0) + predValue.u64 |= 1; + break; + case F32_TYPE: + if(setValue.f32 == 0) + predValue.u64 |= 1; + break; + case F64_TYPE: + if(setValue.f64 == 0) + predValue.u64 |= 1; + break; + default: assert(0); break; + } + + + if ( (type==S8_TYPE)||(type==S16_TYPE)||(type==S32_TYPE)||(type==S64_TYPE)|| + (type==U8_TYPE)||(type==U16_TYPE)||(type==U32_TYPE)||(type==U64_TYPE)|| + (type==B8_TYPE)||(type==B16_TYPE)||(type==B32_TYPE)||(type==B64_TYPE)) { + if((setValue.u32 & (1<<(size-1))) != 0) + predValue.u64 |= 1<<1; + } + if ( type==F32_TYPE ) { + if(setValue.f32 < 0) + predValue.u64 |= 1<<1; + } + + if(dst.get_operand_lohi() == 1) + { + setValue.u64 = ((m_regs.back()[ regName ].u64) & (~(0xFFFF))) + (data.u64 & 0xFFFF); + } + else if(dst.get_operand_lohi() == 2) + { + setValue.u64 = ((m_regs.back()[ regName ].u64) & (~(0xFFFF0000))) + ((data.u64<<16) & 0xFFFF0000); + } + + set_reg(predName,predValue); + set_reg(regName,setValue); + } + else if (type == BB128_TYPE) + { + //b128 stuff here. + ptx_reg_t setValue2, setValue3, setValue4; + setValue.u64 = 0; + setValue2.u64 = 0; + setValue3.u64 = 0; + setValue4.u64 = 0; + setValue.u32 = data.u128.lowest; + setValue2.u32 = data.u128.low; + setValue3.u32 = data.u128.high; + setValue4.u32 = data.u128.highest; + + const symbol *name1, *name2, *name3, *name4 = NULL; + + name1 = dst.vec_symbol(0); + name2 = dst.vec_symbol(1); + name3 = dst.vec_symbol(2); + name4 = dst.vec_symbol(3); + + set_reg(name1,setValue); + set_reg(name2,setValue2); + set_reg(name3,setValue3); + set_reg(name4,setValue4); + } + else if (type == BB64_TYPE) + { + //ptxplus version of storing 64 bit values to registers stores to two adjacent registers + ptx_reg_t setValue2; + setValue.u32 = 0; + setValue2.u32 = 0; + + setValue.u32 = data.bits.ls; + setValue2.u32 = data.bits.ms; + + const symbol *name1, *name2 = NULL; + + name1 = dst.vec_symbol(0); + name2 = dst.vec_symbol(1); + + set_reg(name1,setValue); + set_reg(name2,setValue2); + } + else + { + if(dst.get_operand_lohi() == 1) + { + setValue.u64 = ((m_regs.back()[ dst.get_symbol() ].u64) & (~(0xFFFF))) + (data.u64 & 0xFFFF); + } + else if(dst.get_operand_lohi() == 2) + { + setValue.u64 = ((m_regs.back()[ dst.get_symbol() ].u64) & (~(0xFFFF0000))) + ((data.u64<<16) & 0xFFFF0000); + } + set_reg(dst.get_symbol(),setValue); + } + } + + // global memory - g[4], g[$r0] + else if(dst.get_addr_space() == 1) + { + dstData = thread->get_operand_value(dst, dst, type, thread, 0); + mem = g_global_mem; + type_info_key::type_decode(type,size,t); + + mem->write(dstData.u32,size/8,&data.u128,thread,pI); + thread->m_last_effective_address = dstData.u32; + thread->m_last_memory_space = global_space; + } + + // shared memory - s[4], s[$r0] + else if(dst.get_addr_space() == 2) + { + dstData = thread->get_operand_value(dst, dst, type, thread, 0); + mem = thread->m_shared_mem; + type_info_key::type_decode(type,size,t); + + mem->write(dstData.u32,size/8,&data.u128,thread,pI); + thread->m_last_effective_address = dstData.u32; + thread->m_last_memory_space = shared_space; } - m_last_set_operand_value = data; + + // local memory - l0[4], l0[$r0] + else if(dst.get_addr_space() == 4) + { + dstData = thread->get_operand_value(dst, dst, type, thread, 0); + mem = thread->m_local_mem; + type_info_key::type_decode(type,size,t); + + mem->write(dstData.u32,size/8,&data.u128,thread,pI); + thread->m_last_effective_address = dstData.u32; + thread->m_last_memory_space = local_space; + } + + else + { + printf("Destination stores to unknown location."); + assert(0); + } + + } void ptx_thread_info::set_vector_operand_values( const operand_info &dst, @@ -271,12 +654,12 @@ void ptx_thread_info::set_vector_operand_values( const operand_info &dst, const ptx_reg_t &data4, unsigned num_elements ) { - set_operand_value(dst.vec_symbol(0), data1); - set_operand_value(dst.vec_symbol(1), data2); + set_reg(dst.vec_symbol(0), data1); + set_reg(dst.vec_symbol(1), data2); if (num_elements > 2) { - set_operand_value(dst.vec_symbol(2), data3); + set_reg(dst.vec_symbol(2), data3); if (num_elements > 3) { - set_operand_value(dst.vec_symbol(3), data4); + set_reg(dst.vec_symbol(3), data4); } } @@ -303,13 +686,18 @@ void abs_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case S16_TYPE: d.s16 = my_abs(a.s16); break; case S32_TYPE: d.s32 = my_abs(a.s32); break; case S64_TYPE: d.s64 = my_abs(a.s64); break; + case U16_TYPE: d.s16 = my_abs(a.u16); break; + case U32_TYPE: d.s32 = my_abs(a.u32); break; + case U64_TYPE: d.s64 = my_abs(a.u64); break; case F32_TYPE: d.f32 = my_abs(a.f32); break; case F64_TYPE: d.f64 = my_abs(a.f64); break; default: @@ -318,18 +706,22 @@ void abs_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t src1_data, src2_data, data; + int overflow = 0; + int carry = 0; const operand_info &dst = pI->dst(); //get operand info of sources and destination const operand_info &src1 = pI->src1(); //use them to determine that they are of type 'register' const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); //get values from the operand infos - src2_data = thread->get_operand_value(src2); + + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); unsigned rounding_mode = pI->rounding_mode(); int orig_rm = fegetround(); @@ -339,25 +731,49 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) default: assert(0); break; } - unsigned to_type = pI->get_type(); - - switch ( to_type ) { + //performs addition. Sets carry and overflow if needed. + switch ( i_type ) { case S8_TYPE: + data.s64 = (src1_data.s64 & 0x0000000FF) + (src2_data.s64 & 0x0000000FF); + if(((src1_data.s64 & 0x80)-(src2_data.s64 & 0x80)) == 0) {overflow=((src1_data.s64 & 0x80)-(data.s64 & 0x80))==0?0:1; } + carry = (data.u64 & 0x000000100)>>8; + break; case S16_TYPE: + data.s64 = (src1_data.s64 & 0x00000FFFF) + (src2_data.s64 & 0x00000FFFF); + if(((src1_data.s64 & 0x8000)-(src2_data.s64 & 0x8000)) == 0) {overflow=((src1_data.s64 & 0x8000)-(data.s64 & 0x8000))==0?0:1; } + carry = (data.u64 & 0x000010000)>>16; + break; case S32_TYPE: + data.s64 = (src1_data.s64 & 0x0FFFFFFFF) + (src2_data.s64 & 0x0FFFFFFFF); + if(((src1_data.s64 & 0x80000000)-(src2_data.s64 & 0x80000000)) == 0) {overflow=((src1_data.s64 & 0x80000000)-(data.s64 & 0x80000000))==0?0:1; } + carry = (data.u64 & 0x100000000)>>32; + break; case S64_TYPE: + data.s64 = src1_data.s64 + src2_data.s64; + break; case U8_TYPE: + data.u64 = (src1_data.u64 & 0xFF) + (src2_data.u64 & 0xFF); + carry = (data.u64 & 0x100)>>8; + break; case U16_TYPE: + data.u64 = (src1_data.u64 & 0xFFFF) + (src2_data.u64 & 0xFFFF); + carry = (data.u64 & 0x10000)>>16; + break; case U32_TYPE: + data.u64 = (src1_data.u64 & 0xFFFFFFFF) + (src2_data.u64 & 0xFFFFFFFF); + carry = (data.u64 & 0x100000000)>>32; + break; case U64_TYPE: - data.s64 = src1_data.s64 + src2_data.s64; break; + data.s64 = src1_data.s64 + src2_data.s64; + break; case F16_TYPE: assert(0); break; case F32_TYPE: data.f32 = src1_data.f32 + src2_data.f32; break; case F64_TYPE: data.f64 = src1_data.f64 + src2_data.f64; break; default: assert(0); break; } fesetround( orig_rm ); - thread->set_operand_value(dst,data); + + thread->set_operand_value(dst, data, i_type, thread, pI, overflow, carry ); } void addc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -370,12 +786,45 @@ void and_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); - src2_data = thread->get_operand_value(src2); + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = ~(~(src1_data.pred) & ~(src2_data.pred)); + else + data.u64 = src1_data.u64 & src2_data.u64; + + thread->set_operand_value(dst,data, i_type, thread, pI); +} + +void andn_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t src1_data, src2_data, data; + + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + switch ( i_type ) { + case B16_TYPE: src2_data.u16 = ~src2_data.u16; break; + case B32_TYPE: src2_data.u32 = ~src2_data.u32; break; + case B64_TYPE: src2_data.u64 = ~src2_data.u64; break; + default: + printf("Execution error: type mismatch with instruction\n"); + assert(0); + break; + } data.u64 = src1_data.u64 & src2_data.u64; - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI); } void atom_callback( void* ptx_inst, void* thd ) @@ -406,13 +855,13 @@ void atom_callback( void* ptx_inst, void* thd ) const operand_info &src2 = pI->src2(); // b // Get operand values - src1_data = thread->get_operand_value(src1); // a - src2_data = thread->get_operand_value(src2); // b + src1_data = thread->get_operand_value(src1, dst, to_type, thread, 1); // a + src2_data = thread->get_operand_value(src2, dst, to_type, thread, 1); // b // Copy value pointed to in operand 'a' into register 'd' // (i.e. copy src1_data to dst) g_global_mem->read(src1_data.u32,size/8,&data.s64); - thread->set_operand_value(dst, data); // Write value into register 'd' + thread->set_operand_value(dst, data, to_type, thread, pI); // Write value into register 'd' // Get the atomic operation to be performed unsigned m_atomic_spec = pI->get_atomic(); @@ -490,7 +939,7 @@ void atom_callback( void* ptx_inst, void* thd ) ptx_reg_t src3_data; const operand_info &src3 = pI->src3(); - src3_data = thread->get_operand_value(src3); + src3_data = thread->get_operand_value(src3, dst, to_type, thread, 1); switch ( to_type ) { case B32_TYPE: @@ -648,7 +1097,9 @@ void atom_impl( const ptx_instruction *pI, ptx_thread_info *thread ) // get the memory address const operand_info &src1 = pI->src1(); - ptx_reg_t src1_data = thread->get_operand_value(src1); + const operand_info &dst = pI->dst(); + unsigned i_type = pI->get_type(); + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); memory_space_t space = pI->get_space(); @@ -661,7 +1112,7 @@ void atom_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void bar_sync_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &dst = pI->dst(); - ptx_reg_t b = thread->get_operand_value(dst); + ptx_reg_t b = thread->get_operand_value(dst, dst, U32_TYPE, thread, 1); assert( b.u32 == 0 ); // not clear what should happen if this is not zero } @@ -672,12 +1123,28 @@ void bfind_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not void bra_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &target = pI->dst(); - ptx_reg_t target_pc = thread->get_operand_value(target); + ptx_reg_t target_pc = thread->get_operand_value(target, target, U32_TYPE, thread, 1); thread->m_branch_taken = true; thread->set_npc(target_pc); } +void break_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + const operand_info &target = thread->pop_breakaddr(); + ptx_reg_t target_pc = thread->get_operand_value(target, target, U32_TYPE, thread, 1); + + thread->m_branch_taken = true; + thread->set_npc(target_pc); +} + +void breakaddr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + const operand_info &target = pI->dst(); + thread->push_breakaddr(target); + assert(pI->has_pred() == false); // pdom analysis cannot handle if this instruction is predicated +} + void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void brkpt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -747,21 +1214,22 @@ void cnot_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, b, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + switch ( i_type ) { - case PRED_TYPE: d.pred = (a.pred == 0)?1:0; break; + case PRED_TYPE: d.pred = ((a.pred & 0x0001) == 0)?1:0; break; case B16_TYPE: d.u16 = (a.u16 == 0)?1:0; break; case B32_TYPE: d.u32 = (a.u32 == 0)?1:0; break; case B64_TYPE: d.u64 = (a.u64 == 0)?1:0; break; default: printf("Execution error: type mismatch with instruction\n"); - assert(0); // TODO: add more typechecking like this + assert(0); break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void cos_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -769,9 +1237,11 @@ void cos_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: d.f32 = cos(a.f32); @@ -782,7 +1252,7 @@ void cos_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } ptx_reg_t chop( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, int rounding_mode, int saturation_mode ) @@ -1270,13 +1740,54 @@ void cvt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) unsigned src_fmt = type_info_key::type_decode(from_type, from_width, from_sign); unsigned dst_fmt = type_info_key::type_decode(to_type, to_width, to_sign); - ptx_reg_t data = thread->get_operand_value(src1); + ptx_reg_t data = thread->get_operand_value(src1, dst, from_type, thread, 1); + + if(pI->is_neg()){ + + switch( from_type ) { + // Default to f32 for now, need to add support for others + case S8_TYPE: + case U8_TYPE: + case B8_TYPE: + data.s8 = -data.s8; + break; + case S16_TYPE: + case U16_TYPE: + case B16_TYPE: + data.s16 = -data.s16; + break; + case S32_TYPE: + case U32_TYPE: + case B32_TYPE: + data.s32 = -data.s32; + break; + case S64_TYPE: + case U64_TYPE: + case B64_TYPE: + data.s64 = -data.s64; + break; + case F16_TYPE: + data.f16 = -data.f16; + break; + case F32_TYPE: + data.f32 = -data.f32; + break; + case F64_TYPE: + data.f64 = -data.f64; + break; + default: + assert(0); + } + + } + + if ( g_cvt_fn[src_fmt][dst_fmt] != NULL ) { ptx_reg_t result = g_cvt_fn[src_fmt][dst_fmt](data,from_width,to_width,to_sign, rounding_mode, saturation_mode); data = result; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst, data, to_type, thread, pI ); } void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1288,7 +1799,8 @@ void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread ) memory_space_t space = pI->get_space(); bool to_non_generic = pI->is_to(); - ptx_reg_t from_addr = thread->get_operand_value(src1); + unsigned i_type = pI->get_type(); + ptx_reg_t from_addr = thread->get_operand_value(src1,dst,i_type,thread,1); addr_t from_addr_hw = (addr_t)from_addr.u64; addr_t to_addr_hw = 0; unsigned smid = thread->get_hw_sid(); @@ -1312,7 +1824,7 @@ void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t to_addr; to_addr.u64 = to_addr_hw; - thread->set_operand_value(dst,to_addr); + thread->set_reg(dst.get_symbol(),to_addr); } void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1323,11 +1835,12 @@ void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - ptx_reg_t src1_data = thread->get_operand_value(src1); - ptx_reg_t src2_data = thread->get_operand_value(src2); - unsigned i_type = pI->get_type(); + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + switch ( i_type ) { case S8_TYPE: case S16_TYPE: @@ -1348,7 +1861,7 @@ void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) case F64_TYPE: data.f64 = src1_data.f64 / src2_data.f64; break; default: assert(0); break; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread,pI); } void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1356,9 +1869,12 @@ void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t src1_data, src2_data, data; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - src1_data = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: data.f32 = cuda_math::__powf(2.0, src1_data.f32); @@ -1369,7 +1885,7 @@ void ex2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread,pI); } void exit_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1398,7 +1914,7 @@ void isspacep_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); memory_space_t space = pI->get_space(); - a = thread->get_operand_value(src1); + a = thread->get_reg(src1.get_symbol()); addr_t addr = (addr_t)a.u64; unsigned smid = thread->get_hw_sid(); unsigned hwtid = thread->get_hw_tid(); @@ -1413,7 +1929,7 @@ void isspacep_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t p; p.pred = t?1:0; - thread->set_operand_value(dst,p); + thread->set_reg(dst.get_symbol(),p); } void decode_space( memory_space_t &space, const ptx_thread_info *thread, const operand_info &op, memory_space *&mem, addr_t &addr) @@ -1472,11 +1988,14 @@ void ld_exec( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - ptx_reg_t src1_data = thread->get_operand_value(src1); + + unsigned type = pI->get_type(); + + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, type, thread, 1); ptx_reg_t data; memory_space_t space = pI->get_space(); unsigned vector_spec = pI->get_vector(); - unsigned type = pI->get_type(); + memory_space *mem = NULL; addr_t addr = src1_data.u32; @@ -1490,7 +2009,7 @@ void ld_exec( const ptx_instruction *pI, ptx_thread_info *thread ) mem->read(addr,size/8,&data.s64); if( type == S16_TYPE || type == S32_TYPE ) sign_extend(data,size,dst); - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, type, thread, pI); } else { ptx_reg_t data1, data2, data3, data4; mem->read(addr,size/8,&data1.s64); @@ -1523,9 +2042,12 @@ void lg2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: d.f32 = log(a.f32)/log(2); @@ -1536,7 +2058,7 @@ void lg2_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void mad24_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1546,11 +2068,12 @@ void mad24_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src2 = pI->src2(); const operand_info &src3 = pI->src3(); ptx_reg_t d, t; - ptx_reg_t a = thread->get_operand_value(src1); - ptx_reg_t b = thread->get_operand_value(src2); - ptx_reg_t c = thread->get_operand_value(src3); unsigned i_type = pI->get_type(); + ptx_reg_t a = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t b = thread->get_operand_value(src2, dst, i_type, thread, 1); + ptx_reg_t c = thread->get_operand_value(src3, dst, i_type, thread, 1); + unsigned sat_mode = pI->saturation_mode(); assert( !pI->is_wide() ); @@ -1579,7 +2102,8 @@ void mad24_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert(0); break; } - thread->set_operand_value(dst,d); + + thread->set_operand_value(dst, d, i_type, thread, pI); } void mad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1594,11 +2118,12 @@ void mad_def( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src2 = pI->src2(); const operand_info &src3 = pI->src3(); ptx_reg_t d, t; - ptx_reg_t a = thread->get_operand_value(src1); - ptx_reg_t b = thread->get_operand_value(src2); - ptx_reg_t c = thread->get_operand_value(src3); unsigned i_type = pI->get_type(); + ptx_reg_t a = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t b = thread->get_operand_value(src2, dst, i_type, thread, 1); + ptx_reg_t c = thread->get_operand_value(src3, dst, i_type, thread, 1); + unsigned rounding_mode = pI->rounding_mode(); switch ( i_type ) { @@ -1681,7 +2206,7 @@ void mad_def( const ptx_instruction *pI, ptx_thread_info *thread ) assert(0); break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst, d, i_type, thread, pI); } bool isNaN(float x) @@ -1700,10 +2225,12 @@ void max_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + + switch ( i_type ) { case U16_TYPE: d.u16 = MY_MAX_I(a.u16,b.u16); break; case U32_TYPE: d.u32 = MY_MAX_I(a.u32,b.u32); break; @@ -1719,7 +2246,7 @@ void max_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -1730,10 +2257,12 @@ void min_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + + switch ( i_type ) { case U16_TYPE: d.u16 = MY_MIN_I(a.u16,b.u16); break; case U32_TYPE: d.u32 = MY_MIN_I(a.u32,b.u32); break; @@ -1749,7 +2278,7 @@ void min_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1758,8 +2287,9 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); + unsigned i_type = pI->get_type(); - if( src1.is_vector() || dst.is_vector() ) { + if( (src1.is_vector() || dst.is_vector()) && (i_type != BB64_TYPE) && (i_type != BB128_TYPE) ) { // pack or unpack operation unsigned nbits_to_move; ptx_reg_t tmp_bits; @@ -1786,7 +2316,7 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } } } else { - data = thread->get_operand_value(src1); + data = thread->get_operand_value(src1, dst, i_type, thread, 1); switch( pI->get_type() ) { case B16_TYPE: tmp_bits.u16 = data.u16; break; @@ -1813,11 +2343,14 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } thread->set_vector_operand_values(dst,v[0],v[1],v[2],v[3],nelem); } else { - thread->set_operand_value(dst,tmp_bits); + thread->set_operand_value(dst,tmp_bits, i_type, thread, pI); } } else { - data = thread->get_operand_value(src1); - thread->set_operand_value(dst,data); + + data = thread->get_operand_value(src1, dst, i_type, thread, 1); + + thread->set_operand_value(dst, data, i_type, thread, pI); + } } @@ -1829,13 +2362,17 @@ void mul24_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); - src2_data = thread->get_operand_value(src2); + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + + //src1_data = srcOperandModifiers(src1_data, src1, dst, i_type, thread); + //src2_data = srcOperandModifiers(src2_data, src2, dst, i_type, thread); + src1_data.mask_and(0,0x00FFFFFF); src2_data.mask_and(0,0x00FFFFFF); - unsigned i_type = pI->get_type(); - switch ( i_type ) { case S32_TYPE: if( src1_data.get_bit(23) ) @@ -1860,7 +2397,7 @@ void mul24_impl( const ptx_instruction *pI, ptx_thread_info *thread ) data.mask_and(0,0xFFFFFFFF); } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst, data, i_type, thread, pI); } void mul_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1871,10 +2408,11 @@ void mul_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); ptx_reg_t d, t; - ptx_reg_t a = thread->get_operand_value(src1); - ptx_reg_t b = thread->get_operand_value(src2); unsigned i_type = pI->get_type(); + ptx_reg_t a = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t b = thread->get_operand_value(src2, dst, i_type, thread, 1); + unsigned rounding_mode = pI->rounding_mode(); switch ( i_type ) { @@ -1930,7 +2468,9 @@ void mul_impl( const ptx_instruction *pI, ptx_thread_info *thread ) case RZ_OPTION: fesetround( FE_TOWARDZERO ); break; default: assert(0); break; } + d.f32 = a.f32 * b.f32; + if ( pI->saturation_mode() ) { if ( d.f32 < 0 ) d.f32 = 0; else if ( d.f32 > 1.0f ) d.f32 = 1.0f; @@ -1957,7 +2497,8 @@ void mul_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert(0); break; } - thread->set_operand_value(dst,d); + + thread->set_operand_value(dst, d, i_type, thread, pI); } void neg_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1967,9 +2508,10 @@ void neg_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - src1_data = thread->get_operand_value(src1); - unsigned to_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, to_type, thread, 1); + + switch ( to_type ) { case S8_TYPE: case S16_TYPE: @@ -1987,7 +2529,55 @@ void neg_impl( const ptx_instruction *pI, ptx_thread_info *thread ) default: assert(0); break; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, to_type, thread, pI); +} + +//nandn bitwise negates second operand then bitwise nands with the first operand +void nandn_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t src1_data, src2_data, data; + + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = (~src1_data.pred & src2_data.pred); + else + data.u64 = ~(src1_data.u64 & ~src2_data.u64); + + thread->set_operand_value(dst,data, i_type, thread, pI); + +} + +//norn bitwise negates first operand then bitwise ands with the second operand +void norn_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t src1_data, src2_data, data; + + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = ~(src1_data.pred & ~(src2_data.pred)); + else + data.u64 = ~(src1_data.u64) & src2_data.u64; + + thread->set_operand_value(dst,data, i_type, thread, pI); + } void not_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1995,11 +2585,13 @@ void not_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, b, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { - case PRED_TYPE: d.pred = ~a.pred; break; + case PRED_TYPE: d.pred = (~(a.pred) & 0x000F); break; case B16_TYPE: d.u16 = ~a.u16; break; case B32_TYPE: d.u32 = ~a.u32; break; case B64_TYPE: d.u64 = ~a.u64; break; @@ -2009,7 +2601,7 @@ void not_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void or_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2018,12 +2610,38 @@ void or_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); - src2_data = thread->get_operand_value(src2); - data.u64 = src1_data.u64 | src2_data.u64; + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = ~(~(src1_data.pred) | ~(src2_data.pred)); + else + data.u64 = src1_data.u64 | src2_data.u64; + + thread->set_operand_value(dst,data, i_type, thread, pI); +} + +void orn_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + ptx_reg_t src1_data, src2_data, data; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = ~(~(src1_data.pred) | (src2_data.pred)); + else + data.u64 = src1_data.u64 | ~src2_data.u64; - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI); } void pmevent_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -2037,9 +2655,11 @@ void rcp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t src1_data, src2_data, data; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - src1_data = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: data.f32 = 1.0f / src1_data.f32; @@ -2053,7 +2673,7 @@ void rcp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI); } void red_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -2066,12 +2686,13 @@ void rem_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); - src2_data = thread->get_operand_value(src2); + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); data.u64 = src1_data.u64 % src2_data.u64; - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI); } void ret_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2091,9 +2712,11 @@ void rsqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: if ( a.f32 < 0 ) { @@ -2125,7 +2748,7 @@ void rsqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } #define SAD(d,a,b,c) d = c + ((a<b) ? (b-a) : (a-b)) @@ -2137,11 +2760,13 @@ void sad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src1(); const operand_info &src3 = pI->src1(); - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - c = thread->get_operand_value(src3); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + c = thread->get_operand_value(src3, dst, i_type, thread, 1); + + switch ( i_type ) { case U16_TYPE: SAD(d.u16,a.u16,b.u16,c.u16); break; case U32_TYPE: SAD(d.u32,a.u32,b.u32,c.u32); break; @@ -2157,7 +2782,7 @@ void sad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void selp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2169,13 +2794,16 @@ void selp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, b, c, d; - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - c = thread->get_operand_value(src3); + unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + c = thread->get_operand_value(src3, dst, i_type, thread, 1); - d = (c.pred)?a:b; + //predicate value was changed so the lowest bit being set means the zero flag is set. + //As a result, the value of c.pred must be inverted to get proper behavior + d = (!(c.pred & 0x0001))?a:b; - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, PRED_TYPE, thread, pI); } bool isFloat(int type) @@ -2362,18 +2990,19 @@ void setp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert( pI->get_num_operands() < 4 ); // or need to deal with "c" operand / boolOp - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - unsigned type = pI->get_type(); unsigned cmpop = pI->get_cmpop(); + a = thread->get_operand_value(src1, dst, type, thread, 1); + b = thread->get_operand_value(src2, dst, type, thread, 1); t = CmpOp(type,a,b,cmpop); ptx_reg_t data; - data.pred = (t!=0); - thread->set_operand_value(dst,data); + //the way ptxplus handles the zero flag, 1 = false and 0 = true + data.pred = (t==0); //inverting predicate since ptxplus uses "1" for a set zero flag + + thread->set_operand_value(dst,data, PRED_TYPE, thread, pI); } void set_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2387,12 +3016,30 @@ void set_impl( const ptx_instruction *pI, ptx_thread_info *thread ) assert( pI->get_num_operands() < 4 ); // or need to deal with "c" operand / boolOp - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - unsigned src_type = pI->get_type2(); unsigned cmpop = pI->get_cmpop(); + a = thread->get_operand_value(src1, dst, src_type, thread, 1); + b = thread->get_operand_value(src2, dst, src_type, thread, 1); + + // Take abs of first operand if needed + if(pI->is_abs()) { + switch ( src_type ) { + case S16_TYPE: a.s16 = my_abs(a.s16); break; + case S32_TYPE: a.s32 = my_abs(a.s32); break; + case S64_TYPE: a.s64 = my_abs(a.s64); break; + case U16_TYPE: a.u16 = a.u16; break; + case U32_TYPE: a.u32 = my_abs(a.u32); break; + case U64_TYPE: a.u64 = my_abs(a.u64); break; + case F32_TYPE: a.f32 = my_abs(a.f32); break; + case F64_TYPE: a.f64 = my_abs(a.f64); break; + default: + printf("Execution error: type mismatch with instruction\n"); + assert(0); + break; + } + } + t = CmpOp(src_type,a,b,cmpop); ptx_reg_t data; @@ -2402,7 +3049,7 @@ void set_impl( const ptx_instruction *pI, ptx_thread_info *thread ) data.u32 = (t!=0)?0xFFFFFFFF:0; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst, data, pI->get_type(), thread, pI); } @@ -2412,26 +3059,28 @@ void shl_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - - unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + switch ( i_type ) { case B16_TYPE: + case U16_TYPE: if ( b.u16 >= 16 ) d.u16 = 0; else d.u16 = (unsigned short) ((a.u16 << b.u16) & 0xFFFF); break; case B32_TYPE: + case U32_TYPE: if ( b.u32 >= 32 ) d.u32 = 0; else d.u32 = (unsigned) ((a.u32 << b.u32) & 0xFFFFFFFF); break; case B64_TYPE: + case U64_TYPE: if ( b.u32 >= 64 ) d.u64 = 0; else @@ -2443,7 +3092,7 @@ void shl_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst, d, i_type, thread, pI); } void shr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2452,10 +3101,12 @@ void shr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + + switch ( i_type ) { case U16_TYPE: case B16_TYPE: @@ -2522,7 +3173,7 @@ void shr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void sin_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2530,9 +3181,11 @@ void sin_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: d.f32 = sin(a.f32); @@ -2543,7 +3196,7 @@ void sin_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2555,20 +3208,19 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, b, c, d; - a = thread->get_operand_value(src1); - b = thread->get_operand_value(src2); - c = thread->get_operand_value(src3); - - bool t = false; + unsigned i_type = pI->get_type(); unsigned c_type = pI->get_type2(); + bool t = false; + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + b = thread->get_operand_value(src2, dst, i_type, thread, 1); + c = thread->get_operand_value(src3, dst, c_type, thread, 1); + switch ( c_type ) { case S32_TYPE: t = c.s32 >= 0; break; case F32_TYPE: t = c.f32 >= 0; break; default: assert(0); } - unsigned i_type = pI->get_type(); - switch ( i_type ) { case B16_TYPE: case U16_TYPE: d.u16 = t?a.u16:b.u16; break; @@ -2581,7 +3233,7 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread ) default: assert(0); } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -2589,9 +3241,11 @@ void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t a, d; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); - a = thread->get_operand_value(src1); unsigned i_type = pI->get_type(); + a = thread->get_operand_value(src1, dst, i_type, thread, 1); + + switch ( i_type ) { case F32_TYPE: if ( a.f32 < 0 ) @@ -2609,18 +3263,19 @@ void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) break; } - thread->set_operand_value(dst,d); + thread->set_operand_value(dst,d, i_type, thread, pI); } void st_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); //may be scalar or vector of regs - ptx_reg_t addr_reg = thread->get_operand_value(dst); + unsigned type = pI->get_type(); + ptx_reg_t addr_reg = thread->get_operand_value(dst, dst, type, thread, 1); ptx_reg_t data; memory_space_t space = pI->get_space(); unsigned vector_spec = pI->get_vector(); - unsigned type = pI->get_type(); + memory_space *mem = NULL; addr_t addr = addr_reg.u32; @@ -2631,7 +3286,7 @@ void st_impl( const ptx_instruction *pI, ptx_thread_info *thread ) type_info_key::type_decode(type,size,t); if (!vector_spec) { - data = thread->get_operand_value(src1); + data = thread->get_operand_value(src1, dst, type, thread, 1); mem->write(addr,size/8,&data.s64,thread,pI); } else { if (vector_spec == V2_TYPE) { @@ -2666,30 +3321,54 @@ void st_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void sub_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t data; + int overflow = 0; + int carry = 0; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - ptx_reg_t src1_data = thread->get_operand_value(src1); - ptx_reg_t src2_data = thread->get_operand_value(src2); - unsigned i_type = pI->get_type(); + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + //performs addition. Sets carry and overflow if needed. + //the constant is added in during subtraction so the carry bit is set properly. switch ( i_type ) { case S8_TYPE: + data.s64 = (src1_data.s64 & 0xFF) - (src2_data.s64 & 0xFF) + 0x100; + if(((src1_data.s64 & 0x80)-(src2_data.s64 & 0x80)) != 0) {overflow=((src1_data.s64 & 0x80)-(data.s64 & 0x80))==0?0:1; } + carry = (data.s32 & 0x100)>>8; + break; case S16_TYPE: + data.s64 = (src1_data.s64 & 0xFFFF) - (src2_data.s64 & 0xFFFF) + 0x10000; + if(((src1_data.s64 & 0x8000)-(src2_data.s64 & 0x8000)) != 0) {overflow=((src1_data.s64 & 0x8000)-(data.s64 & 0x8000))==0?0:1; } + carry = (data.s32 & 0x10000)>>16; + break; case S32_TYPE: + data.s64 = (src1_data.s64 & 0xFFFFFFFF) - (src2_data.s64 & 0xFFFFFFFF) + 0x100000000; + if(((src1_data.s64 & 0x80000000)-(src2_data.s64 & 0x80000000)) != 0) {overflow=((src1_data.s64 & 0x80000000)-(data.s64 & 0x80000000))==0?0:1; } + carry = ((data.u64)>>32) & 0x0001; + break; case S64_TYPE: data.s64 = src1_data.s64 - src2_data.s64; break; - case U8_TYPE: - case U16_TYPE: - case U32_TYPE: - case U64_TYPE: case B8_TYPE: + case U8_TYPE: + data.u64 = (src1_data.u64 & 0xFF) - (src2_data.u64 & 0xFF) + 0x100; + carry = (data.u64 & 0x100)>>8; + break; case B16_TYPE: + case U16_TYPE: + data.u64 = (src1_data.u64 & 0xFFFF) - (src2_data.u64 & 0xFFFF) + 0x10000; + carry = (data.u64 & 0x10000)>>16; + break; case B32_TYPE: + case U32_TYPE: + data.u64 = (src1_data.u64 & 0xFFFFFFFF) - (src2_data.u64 & 0xFFFFFFFF) + 0x100000000; + carry = (data.u64 & 0x100000000)>>32; + break; case B64_TYPE: + case U64_TYPE: data.u64 = src1_data.u64 - src2_data.u64; break; case F16_TYPE: assert(0); break; case F32_TYPE: data.f32 = src1_data.f32 - src2_data.f32; break; @@ -2697,7 +3376,12 @@ void sub_impl( const ptx_instruction *pI, ptx_thread_info *thread ) default: assert(0); break; } - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI, overflow, carry); +} + +void nop_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + // Do nothing } void subc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } @@ -2777,7 +3461,8 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) fflush(stdout); ptx_reg_t data1, data2, data3, data4; if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4]; - thread->get_vector_operand_values(src2, ptx_tex_regs, 4); //ptx_reg should be 4 entry vector type...coordinates into texture + unsigned nelem = src2.get_vect_nelem(); + thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture assert(NameToTextureMap.find(texname) != NameToTextureMap.end());//use map to find texturerefence, then use map to find pointer to array struct textureReference* texref = NameToTextureMap[texname]; @@ -2905,6 +3590,9 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) case U8_TYPE: case U16_TYPE: case U32_TYPE: + case B8_TYPE: + case B16_TYPE: + case B32_TYPE: case S8_TYPE: case S16_TYPE: case S32_TYPE: @@ -2918,6 +3606,7 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } } break; + case B64_TYPE: case U64_TYPE: case S64_TYPE: mem->read( tex_array_index, 8, &data1.u64); @@ -3041,16 +3730,17 @@ void vote_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ptx_reg_t src1_data; const operand_info &src1 = pI->src1(); - src1_data = thread->get_operand_value(src1); + src1_data = thread->get_operand_value(src1, pI->dst(), PRED_TYPE, thread, 1); - bool pred_value = src1_data.pred; + //predicate value was changed so the lowest bit being set means the zero flag is set. + //As a result, the value of src1_data.pred must be inverted to get proper behavior + bool pred_value = !(src1_data.pred & 0x0001); bool invert = src1.is_neg_pred(); threads_in_warp.push_back(thread); and_all &= (invert ^ pred_value); or_all |= (invert ^ pred_value); - // TODO: determine last active thread in warp... if( thread->get_hw_tid() == last_tid ) { bool pred_value = false; @@ -3062,11 +3752,11 @@ void vote_impl( const ptx_instruction *pI, ptx_thread_info *thread ) abort(); } ptx_reg_t data; - data.pred = pred_value?1:0; + data.pred = pred_value?0:1; //the way ptxplus handles the zero flag, 1 = false and 0 = true for( std::list<ptx_thread_info*>::iterator t=threads_in_warp.begin(); t!=threads_in_warp.end(); ++t ) { const operand_info &dst = pI->dst(); - (*t)->set_operand_value(dst,data); + (*t)->set_operand_value(dst,data, PRED_TYPE, (*t), pI); } first_in_warp = true; } @@ -3080,12 +3770,17 @@ void xor_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); - src1_data = thread->get_operand_value(src1); - src2_data = thread->get_operand_value(src2); + unsigned i_type = pI->get_type(); + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); - data.u64 = src1_data.u64 ^ src2_data.u64; + //the way ptxplus handles predicates: 1 = false and 0 = true + if(i_type == PRED_TYPE) + data.pred = ~(~(src1_data.pred) ^ ~(src2_data.pred)); + else + data.u64 = src1_data.u64 ^ src2_data.u64; - thread->set_operand_value(dst,data); + thread->set_operand_value(dst,data, i_type, thread, pI); } void inst_not_implemented( const ptx_instruction * pI ) @@ -3101,3 +3796,61 @@ void print_instruction(const ptx_instruction *instruction) { } +ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread) +{ + ptx_reg_t result; + memory_space *mem = NULL; + size_t size; + int t; + result.u64=0; + + //complete other cases for reading from memory, such as reading from other const memory + if(opInfo.get_addr_space() == 1) + { + mem = g_global_mem; + type_info_key::type_decode(type,size,t); + mem->read(opData.u32,size/8,&result.u64); + if( type == S16_TYPE || type == S32_TYPE ) + sign_extend(result,size,dstInfo); + } + else if(opInfo.get_addr_space() == 2) + { + mem = thread->m_shared_mem; + type_info_key::type_decode(type,size,t); + mem->read(opData.u32,size/8,&result.u64); + + if( type == S16_TYPE || type == S32_TYPE ) + sign_extend(result,size,dstInfo); + + } + else if(opInfo.get_addr_space() == 3) + { + mem = g_global_mem; + type_info_key::type_decode(type,size,t); + + mem->read((opData.u32 + opInfo.get_const_mem_offset()),size/8,&result.u64); + + if( type == S16_TYPE || type == S32_TYPE ) + sign_extend(result,size,dstInfo); + } + else + { + result = opData; + } + + if(opInfo.get_operand_lohi() == 1) + { + result.u64 = result.u64 & 0xFFFF; + } + else if(opInfo.get_operand_lohi() == 2) + { + result.u64 = (result.u64>>16) & 0xFFFF; + } + + if(opInfo.get_operand_neg() == true) { + result.f32 = -result.f32; + } + + return result; +} + diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 4566468..c1217ee 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -69,12 +69,14 @@ Control 3 SFU 4 Mem(except Tex) 5 Tex 6 +Nop 7 Other 10 */ OP_DEF(ABS_OP,abs_impl,"abs",1,1) OP_DEF(ADD_OP,add_impl,"add",1,1) OP_DEF(ADDC_OP,addc_impl,"addc",1,1) OP_DEF(AND_OP,and_impl,"and",1,1) +OP_DEF(ANDN_OP,andn_impl,"andn",1,1) OP_DEF(ATOM_OP,atom_impl,"atom",0,3) OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",1,3) OP_DEF(BFE_OP,bfe_impl,"bfe",1,1) @@ -106,8 +108,11 @@ OP_DEF(MOV_OP,mov_impl,"mov",1,1) OP_DEF(MUL24_OP,mul24_impl,"mul24",1,1) OP_DEF(MUL_OP,mul_impl,"mul",1,1) OP_DEF(NEG_OP,neg_impl,"neg",1,1) +OP_DEF(NANDN_OP,nandn_impl,"nandn",1,1) +OP_DEF(NORN_OP,norn_impl,"norn",1,1) OP_DEF(NOT_OP,not_impl,"not",1,1) OP_DEF(OR_OP,or_impl,"or",1,1) +OP_DEF(ORN_OP,orn_impl,"orn",1,1) OP_DEF(PMEVENT_OP,pmevent_impl,"pmevent",1,10) OP_DEF(POPC_OP,popc_impl,"popc",1,1) OP_DEF(PREFETCH_OP,prefetch_impl,"prefetch",1,5) @@ -147,3 +152,6 @@ OP_DEF(VSHR_OP,vshr_impl,"vshr",0,11) OP_DEF(VSUB_OP,vsub_impl,"vsub",0,11) OP_DEF(VOTE_OP,vote_impl,"vote",0,3) OP_DEF(XOR_OP,xor_impl,"xor",1,1) +OP_DEF(NOP_OP,nop_impl,"nop",0,7) +OP_DEF(BREAK_OP,break_impl,"break",0,3) +OP_DEF(BREAKADDR_OP,breakaddr_impl,"breakaddr",0,3) diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index 1a06415..2ceafaf 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -72,6 +72,7 @@ enum opcode_t { enum special_regs { CLOCK_REG, + HALFCLOCK_ID, CLOCK64_REG, CTAID_REG, ENVREG_REG, diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 3b83e44..82daef9 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -88,6 +88,7 @@ abs TC; ptx_lval.int_value = ABS_OP; return OPCODE; add TC; ptx_lval.int_value = ADD_OP; return OPCODE; addc TC; ptx_lval.int_value = ADDC_OP; return OPCODE; and TC; ptx_lval.int_value = AND_OP; return OPCODE; +andn TC; ptx_lval.int_value = ANDN_OP; return OPCODE; atom TC; ptx_lval.int_value = ATOM_OP; return OPCODE; bar.sync TC; ptx_lval.int_value = BAR_OP; return OPCODE; bfe TC; ptx_lval.int_value = BFE_OP; return OPCODE; @@ -120,8 +121,11 @@ mov TC; ptx_lval.int_value = MOV_OP; return OPCODE; mul24 TC; ptx_lval.int_value = MUL24_OP; return OPCODE; mul TC; ptx_lval.int_value = MUL_OP; return OPCODE; neg TC; ptx_lval.int_value = NEG_OP; return OPCODE; +nandn TC; ptx_lval.int_value = NANDN_OP; return OPCODE; +norn TC; ptx_lval.int_value = NORN_OP; return OPCODE; not TC; ptx_lval.int_value = NOT_OP; return OPCODE; or TC; ptx_lval.int_value = OR_OP; return OPCODE; +orn TC; ptx_lval.int_value = ORN_OP; return OPCODE; pmevent TC; ptx_lval.int_value = PMEVENT_OP; return OPCODE; popc TC; ptx_lval.int_value = POPC_OP; return OPCODE; prefetch TC; ptx_lval.int_value = PREFETCH_OP; return OPCODE; @@ -163,6 +167,9 @@ vshr TC; ptx_lval.int_value = VSHR_OP; return OPCODE; vsub TC; ptx_lval.int_value = VSUB_OP; return OPCODE; vote TC; ptx_lval.int_value = VOTE_OP; return OPCODE; xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; +nop TC; ptx_lval.int_value = NOP_OP; return OPCODE; +break TC; ptx_lval.int_value = BREAK_OP; return OPCODE; +breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; <INITIAL,NOT_OPCODE>{ @@ -199,8 +206,11 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.version TC; return VERSION_DIRECTIVE; \.visible TC; return VISIBLE_DIRECTIVE; +\.constptr TC; return CONSTPTR_DIRECTIVE; /* Ptx plus directive for pointer to constant memory */ + "%clock" TC; ptx_lval.int_value = CLOCK_REG; return SPECIAL_REGISTER; +"%halfclock" TC; ptx_lval.int_value = HALFCLOCK_ID; return SPECIAL_REGISTER; "%clock64" TC; ptx_lval.int_value = CLOCK64_REG; return SPECIAL_REGISTER; "%ctaid" TC; ptx_lval.int_value = CTAID_REG; return SPECIAL_REGISTER; "%envreg"[0-9]+ TC; sscanf(yytext+7,"%u",&ptx_lval.int_value); ptx_lval.int_value<<=16; ptx_lval.int_value += ENVREG_REG; return SPECIAL_REGISTER; @@ -249,6 +259,8 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.b16 TC; return B16_TYPE; \.b32 TC; return B32_TYPE; \.b64 TC; return B64_TYPE; +\.bb64 TC; return BB64_TYPE; +\.bb128 TC; return BB128_TYPE; \.pred TC; return PRED_TYPE; \.texref TC; return TEXREF_TYPE; @@ -276,6 +288,9 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.le TC; return LE_OPTION; \.gt TC; return GT_OPTION; \.ge TC; return GE_OPTION; +\.cf TC; return CF_OPTION; +\.sf TC; return SF_OPTION; +\.nsf TC; return NSF_OPTION; \.lo TC; return LO_OPTION; \.ls TC; return LS_OPTION; @@ -295,6 +310,8 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.ftz TC; return FTZ_OPTION; +\.neg TC; return NEG_OPTION; + \.wide TC; return WIDE_OPTION; \.uni TC; return UNI_OPTION; @@ -306,6 +323,10 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.gl TC; return GLOBAL_OPTION; \.cta TC; return CTA_OPTION; +\.exit TC; return EXIT_OPTION; + +\.abs TC; return ABS_OPTION; + \.to TC; return TO_OPTION; \.ca TC; return CA_OPTION; @@ -336,9 +357,11 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; \.y TC; ptx_lval.int_value = 1; return DIMENSION_MODIFIER; \.z TC; ptx_lval.int_value = 2; return DIMENSION_MODIFIER; +"-" TC; return MINUS; "+" TC; return PLUS; "," TC; return COMMA; "@" TC; return PRED; +"|" TC; return PIPE; "[" TC; return LEFT_SQUARE_BRACKET; "]" TC; return RIGHT_SQUARE_BRACKET; "<" TC; return LEFT_ANGLE_BRACKET; @@ -352,6 +375,7 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; "{" TC; BEGIN(INITIAL); return LEFT_BRACE; "}" TC; return RIGHT_BRACE; \. TC; return PERIOD; +"/" TC; return BACKSLASH; "//"[^\n]* TC; // eat single diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index a1e1428..1551f22 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -77,6 +77,7 @@ %token CALLPROTOTYPE_DIRECTIVE %token CALLTARGETS_DIRECTIVE %token <int_value> CONST_DIRECTIVE +%token CONSTPTR_DIRECTIVE %token ENTRY_DIRECTIVE %token EXTERN_DIRECTIVE %token FILE_DIRECTIVE @@ -121,6 +122,8 @@ %token B16_TYPE %token B32_TYPE %token B64_TYPE +%token BB64_TYPE +%token BB128_TYPE %token PRED_TYPE %token TEXREF_TYPE %token SAMPLERREF_TYPE @@ -148,18 +151,24 @@ %token GEU_OPTION %token NUM_OPTION %token NAN_OPTION +%token CF_OPTION +%token SF_OPTION +%token NSF_OPTION %token LEFT_SQUARE_BRACKET %token RIGHT_SQUARE_BRACKET %token WIDE_OPTION %token <int_value> SPECIAL_REGISTER +%token MINUS %token PLUS %token COLON %token SEMI_COLON %token EXCLAMATION +%token PIPE %token RIGHT_BRACE %token LEFT_BRACE %token EQUALS %token PERIOD +%token BACKSLASH %token <int_value> DIMENSION_MODIFIER %token RN_OPTION %token RZ_OPTION @@ -175,6 +184,7 @@ %token GEOM_MODIFIER_3D %token SAT_OPTION %token FTZ_OPTION +%token NEG_OPTION %token ATOMIC_AND %token ATOMIC_OR %token ATOMIC_XOR @@ -195,6 +205,8 @@ %token ALL_OPTION %token GLOBAL_OPTION %token CTA_OPTION +%token EXIT_OPTION +%token ABS_OPTION %token TO_OPTION %token CA_OPTION; %token CG_OPTION; @@ -224,30 +236,32 @@ input: /* empty */ | input function_decl ; -function_defn: function_decl { set_symtab($1); } statement_block { end_function(); } - | function_decl { set_symtab($1); } block_spec statement_block { end_function(); } +function_defn: function_decl { set_symtab($1); func_header(".skip"); } statement_block { end_function(); } + | function_decl { set_symtab($1); } block_spec { func_header(".skip"); } statement_block { end_function(); } ; -block_spec: MAXNTID_DIRECTIVE INT_OPERAND COMMA INT_OPERAND COMMA INT_OPERAND +block_spec: MAXNTID_DIRECTIVE INT_OPERAND COMMA INT_OPERAND COMMA INT_OPERAND {func_header_info_int(".maxntid", $2); + func_header_info_int(",", $4); + func_header_info_int(",", $6); } ; -function_decl: function_decl_header LEFT_PAREN { start_function($1); } param_entry RIGHT_PAREN function_ident_param { $$ = reset_symtab(); } +function_decl: function_decl_header LEFT_PAREN { start_function($1); func_header_info("(");} param_entry RIGHT_PAREN {func_header_info(")");} function_ident_param { $$ = reset_symtab(); } | function_decl_header { start_function($1); } function_ident_param { $$ = reset_symtab(); } | function_decl_header { start_function($1); add_function_name(""); g_func_decl=0; $$ = reset_symtab(); } ; -function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN param_list RIGHT_PAREN { g_func_decl=0; } +function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN {func_header_info("(");} param_list RIGHT_PAREN { g_func_decl=0; func_header_info(")"); } | IDENTIFIER { add_function_name($1); g_func_decl=0; } ; -function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; } - | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; } - | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; } - | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; } +function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".entry"); } + | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; func_header(".func"); } + | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); } ; param_list: param_entry { add_directive(); } - | param_list COMMA param_entry { add_directive(); } + | param_list COMMA {func_header_info(",");} param_entry { add_directive(); } param_entry: PARAM_DIRECTIVE { add_space_spec(param_space_unclassified,0); } variable_spec identifier_spec { add_function_arg(); } | REG_DIRECTIVE { add_space_spec(reg_space,0); } variable_spec identifier_spec { add_function_arg(); } @@ -264,8 +278,8 @@ statement_list: directive_statement { add_directive(); } directive_statement: variable_declaration SEMI_COLON | VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2); } - | TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER - | TARGET_DIRECTIVE IDENTIFIER + | TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER { target_header2($2,$4); } + | TARGET_DIRECTIVE IDENTIFIER { target_header($2); } | FILE_DIRECTIVE INT_OPERAND STRING { add_file($2,$3); } | LOC_DIRECTIVE INT_OPERAND INT_OPERAND INT_OPERAND | PRAGMA_DIRECTIVE STRING SEMI_COLON { add_pragma($2); } @@ -274,6 +288,7 @@ directive_statement: variable_declaration SEMI_COLON variable_declaration: variable_spec identifier_list { add_variables(); } | variable_spec identifier_spec EQUALS initializer_list { add_variables(); } | variable_spec identifier_spec EQUALS literal_operand { add_variables(); } + | CONSTPTR_DIRECTIVE IDENTIFIER COMMA IDENTIFIER COMMA INT_OPERAND { add_constptr($2, $4, $6); } ; variable_spec: var_spec_list { set_variable_type(); } @@ -281,21 +296,21 @@ variable_spec: var_spec_list { set_variable_type(); } identifier_list: identifier_spec | identifier_list COMMA identifier_spec; -identifier_spec: IDENTIFIER { add_identifier($1,0,NON_ARRAY_IDENTIFIER); } - | IDENTIFIER LEFT_ANGLE_BRACKET INT_OPERAND RIGHT_ANGLE_BRACKET { +identifier_spec: IDENTIFIER { add_identifier($1,0,NON_ARRAY_IDENTIFIER); func_header_info($1);} + | IDENTIFIER LEFT_ANGLE_BRACKET INT_OPERAND RIGHT_ANGLE_BRACKET { func_header_info($1); func_header_info_int("<", $3); func_header_info(">"); int i,lbase,l; char *id = NULL; lbase = strlen($1); for( i=0; i < $3; i++ ) { l = lbase + (int)log10(i+1)+10; - id = malloc(l); + id = (char*) malloc(l); snprintf(id,l,"%s%u",$1,i); add_identifier(id,0,NON_ARRAY_IDENTIFIER); } free($1); } - | IDENTIFIER LEFT_SQUARE_BRACKET RIGHT_SQUARE_BRACKET { add_identifier($1,0,ARRAY_IDENTIFIER_NO_DIM); } - | IDENTIFIER LEFT_SQUARE_BRACKET INT_OPERAND RIGHT_SQUARE_BRACKET { add_identifier($1,$3,ARRAY_IDENTIFIER); } + | IDENTIFIER LEFT_SQUARE_BRACKET RIGHT_SQUARE_BRACKET { add_identifier($1,0,ARRAY_IDENTIFIER_NO_DIM); func_header_info($1); func_header_info("["); func_header_info("]");} + | IDENTIFIER LEFT_SQUARE_BRACKET INT_OPERAND RIGHT_SQUARE_BRACKET { add_identifier($1,$3,ARRAY_IDENTIFIER); func_header_info($1); func_header_info_int("[",$3); func_header_info("]");} ; var_spec_list: var_spec @@ -327,16 +342,16 @@ type_spec: scalar_type | vector_spec scalar_type ; -vector_spec: V2_TYPE { add_option(V2_TYPE); } - | V3_TYPE { add_option(V3_TYPE); } - | V4_TYPE { add_option(V4_TYPE); } +vector_spec: V2_TYPE { add_option(V2_TYPE); func_header_info(".v2");} + | V3_TYPE { add_option(V3_TYPE); func_header_info(".v3");} + | V4_TYPE { add_option(V4_TYPE); func_header_info(".v4");} ; -scalar_type: S8_TYPE { add_scalar_type_spec( S8_TYPE ); } +scalar_type: S8_TYPE { add_scalar_type_spec( S8_TYPE ); } | S16_TYPE { add_scalar_type_spec( S16_TYPE ); } | S32_TYPE { add_scalar_type_spec( S32_TYPE ); } | S64_TYPE { add_scalar_type_spec( S64_TYPE ); } - | U8_TYPE { add_scalar_type_spec( U8_TYPE ); } + | U8_TYPE { add_scalar_type_spec( U8_TYPE ); } | U16_TYPE { add_scalar_type_spec( U16_TYPE ); } | U32_TYPE { add_scalar_type_spec( U32_TYPE ); } | U64_TYPE { add_scalar_type_spec( U64_TYPE ); } @@ -347,6 +362,8 @@ scalar_type: S8_TYPE { add_scalar_type_spec( S8_TYPE ); } | B16_TYPE { add_scalar_type_spec( B16_TYPE ); } | B32_TYPE { add_scalar_type_spec( B32_TYPE ); } | B64_TYPE { add_scalar_type_spec( B64_TYPE ); } + | BB64_TYPE { add_scalar_type_spec( BB64_TYPE ); } + | BB128_TYPE { add_scalar_type_spec( BB128_TYPE ); } | PRED_TYPE { add_scalar_type_spec( PRED_TYPE ); } | TEXREF_TYPE { add_scalar_type_spec( TEXREF_TYPE ); } | SAMPLERREF_TYPE { add_scalar_type_spec( SAMPLERREF_TYPE ); } @@ -357,7 +374,7 @@ initializer_list: LEFT_BRACE literal_list RIGHT_BRACE { add_array_initializer(); | LEFT_BRACE initializer_list RIGHT_BRACE { syntax_not_implemented(); } literal_list: literal_operand - | literal_operand COMMA literal_list; + | literal_list COMMA literal_operand; instruction_statement: instruction SEMI_COLON | IDENTIFIER COLON { add_label($1); } @@ -372,8 +389,18 @@ instruction: opcode_spec LEFT_PAREN operand RIGHT_PAREN { set_return(); } COMMA opcode_spec: OPCODE { add_opcode($1); } option_list | OPCODE { add_opcode($1); } -pred_spec: PRED IDENTIFIER { add_pred($2,0); } - | PRED EXCLAMATION IDENTIFIER { add_pred($3,1); } +pred_spec: PRED IDENTIFIER { add_pred($2,0, -1); } + | PRED EXCLAMATION IDENTIFIER { add_pred($3,1, -1); } + | PRED IDENTIFIER EQ_OPTION { add_pred($2,0,2); } + | PRED IDENTIFIER LE_OPTION { add_pred($2,0,3); } + | PRED IDENTIFIER NE_OPTION { add_pred($2,0,5); } + | PRED IDENTIFIER GE_OPTION { add_pred($2,0,6); } + | PRED IDENTIFIER EQU_OPTION { add_pred($2,0,10); } + | PRED IDENTIFIER GTU_OPTION { add_pred($2,0,12); } + | PRED IDENTIFIER NEU_OPTION { add_pred($2,0,13); } + | PRED IDENTIFIER CF_OPTION { add_pred($2,0,17); } + | PRED IDENTIFIER SF_OPTION { add_pred($2,0,19); } + | PRED IDENTIFIER NSF_OPTION { add_pred($2,0,28); } ; option_list: option @@ -394,9 +421,12 @@ option: type_spec | GEOM_MODIFIER_3D { add_option(GEOM_MODIFIER_3D); } | SAT_OPTION { add_option(SAT_OPTION); } | FTZ_OPTION { add_option(FTZ_OPTION); } + | NEG_OPTION { add_option(NEG_OPTION); } | APPROX_OPTION { add_option(APPROX_OPTION); } | FULL_OPTION { add_option(FULL_OPTION); } - | atomic_operation_spec + | EXIT_OPTION { add_option(EXIT_OPTION); } + | ABS_OPTION { add_option(ABS_OPTION); } + | atomic_operation_spec ; | TO_OPTION { add_option(TO_OPTION); } ; @@ -452,31 +482,51 @@ operand_list: operand operand: IDENTIFIER { add_scalar_operand( $1 ); } | EXCLAMATION IDENTIFIER { add_neg_pred_operand( $2 ); } + | MINUS IDENTIFIER { add_scalar_operand( $2 ); change_operand_neg(); } | memory_operand | literal_operand | builtin_operand | vector_operand | tex_operand | IDENTIFIER PLUS INT_OPERAND { add_address_operand($1,$3); } + | IDENTIFIER LO_OPTION { add_scalar_operand( $1 ); change_operand_lohi(1);} + | IDENTIFIER HI_OPTION { add_scalar_operand( $1 ); change_operand_lohi(2);} + | IDENTIFIER PIPE IDENTIFIER { add_2vector_operand($1,$3); change_double_operand_type(-1);} + | IDENTIFIER PIPE IDENTIFIER LO_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(1);} + | IDENTIFIER PIPE IDENTIFIER HI_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-1); change_operand_lohi(2);} + | IDENTIFIER BACKSLASH IDENTIFIER { add_2vector_operand($1,$3); change_double_operand_type(-3);} + | IDENTIFIER BACKSLASH IDENTIFIER LO_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-3); change_operand_lohi(1);} + | IDENTIFIER BACKSLASH IDENTIFIER HI_OPTION { add_2vector_operand($1,$3); change_double_operand_type(-3); change_operand_lohi(2);} ; vector_operand: LEFT_BRACE IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_2vector_operand($2,$4); } - | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_3vector_operand($2,$4,$6); } - | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_4vector_operand($2,$4,$6,$8); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_3vector_operand($2,$4,$6); } + | LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE { add_4vector_operand($2,$4,$6,$8); } ; -tex_operand: LEFT_SQUARE_BRACKET IDENTIFIER COMMA - LEFT_BRACE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER RIGHT_BRACE - RIGHT_SQUARE_BRACKET { - add_scalar_operand($2); - add_4vector_operand($5,$7,$9,$11); - } +tex_operand: LEFT_SQUARE_BRACKET IDENTIFIER COMMA { add_scalar_operand($2); } + vector_operand + RIGHT_SQUARE_BRACKET builtin_operand: SPECIAL_REGISTER DIMENSION_MODIFIER { add_builtin_operand($1,$2); } | SPECIAL_REGISTER { add_builtin_operand($1,-1); } ; memory_operand : LEFT_SQUARE_BRACKET address_expression RIGHT_SQUARE_BRACKET { add_memory_operand(); } + | IDENTIFIER LEFT_SQUARE_BRACKET address_expression RIGHT_SQUARE_BRACKET { add_memory_operand(); change_memory_addr_space($1); } + | IDENTIFIER LEFT_SQUARE_BRACKET literal_operand RIGHT_SQUARE_BRACKET { change_memory_addr_space($1); } + | IDENTIFIER LEFT_SQUARE_BRACKET twin_operand RIGHT_SQUARE_BRACKET { change_memory_addr_space($1); add_memory_operand();} + | MINUS memory_operand { change_operand_neg(); } + ; + +twin_operand : IDENTIFIER PLUS IDENTIFIER { add_double_operand($1,$3); change_double_operand_type(1); } + | IDENTIFIER PLUS IDENTIFIER LO_OPTION { add_double_operand($1,$3); change_double_operand_type(1); change_operand_lohi(1); } + | IDENTIFIER PLUS IDENTIFIER HI_OPTION { add_double_operand($1,$3); change_double_operand_type(1); change_operand_lohi(2); } + | IDENTIFIER PLUS EQUALS IDENTIFIER { add_double_operand($1,$4); change_double_operand_type(2); } + | IDENTIFIER PLUS EQUALS IDENTIFIER LO_OPTION { add_double_operand($1,$4); change_double_operand_type(2); change_operand_lohi(1); } + | IDENTIFIER PLUS EQUALS IDENTIFIER HI_OPTION { add_double_operand($1,$4); change_double_operand_type(2); change_operand_lohi(2); } + | IDENTIFIER PLUS EQUALS INT_OPERAND { add_address_operand($1,$4); change_double_operand_type(3); } + ; literal_operand : INT_OPERAND { add_literal_int($1); } | FLOAT_OPERAND { add_literal_float($1); } @@ -484,6 +534,8 @@ literal_operand : INT_OPERAND { add_literal_int($1); } ; address_expression: IDENTIFIER { add_address_operand($1,0); } + | IDENTIFIER LO_OPTION { add_address_operand($1,0); change_operand_lohi(1);} + | IDENTIFIER HI_OPTION { add_address_operand($1,0); change_operand_lohi(2); } | IDENTIFIER PLUS INT_OPERAND { add_address_operand($1,$3); } ; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 5945b72..cd925a5 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -459,7 +459,53 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f assert(pI->get_opcode() == BRA_OP); } } +bool function_info::connect_break_targets() //connecting break instructions with proper targets +{ + std::vector<basic_block_t*>::iterator bb_itr; + std::vector<basic_block_t*>::iterator bb_target_itr; + bool modified = false; + + //start from first basic block, which we know is the entry point + bb_itr = m_basic_blocks.begin(); + for (bb_itr = m_basic_blocks.begin();bb_itr != m_basic_blocks.end(); bb_itr++) { + basic_block_t *p_bb = *bb_itr; + ptx_instruction *pI = p_bb->ptx_end; + if (p_bb->is_exit) //reached last basic block, no successors to link + continue; + if (pI->get_opcode() == BREAK_OP) { + // backup existing successor_ids for stability check + std::set<int> orig_successor_ids = p_bb->successor_ids; + + // erase the previous linkage with old successors + for(std::set<int>::iterator succ_ids = p_bb->successor_ids.begin(); succ_ids != p_bb->successor_ids.end(); ++succ_ids) { + basic_block_t *successor_bb = m_basic_blocks[*succ_ids]; + successor_bb->predecessor_ids.erase(p_bb->bb_id); + } + p_bb->successor_ids.clear(); + //find successor and link that basic_block to this one + //successor of a break is set by an preceeding breakaddr instruction + operand_info *target = find_break_target(pI); + unsigned addr = labels[ target->name() ]; + ptx_instruction *target_pI = m_instr_mem[addr]; + basic_block_t *target_bb = target_pI->get_bb(); + p_bb->successor_ids.insert(target_bb->bb_id); + target_bb->predecessor_ids.insert(p_bb->bb_id); + + if (pI->has_pred()) { + // predicated break - add link to next basic block + unsigned next_addr = pI->get_m_instr_mem_index() + 1; + basic_block_t *next_bb = m_instr_mem[next_addr]->get_bb(); + p_bb->successor_ids.insert(next_bb->bb_id); + next_bb->predecessor_ids.insert(p_bb->bb_id); + } + + modified = modified || (orig_successor_ids != p_bb->successor_ids); + } + } + + return modified; +} void intersect( std::set<int> &A, const std::set<int> &B ) { // return intersection of A and B in A @@ -493,6 +539,44 @@ void print_set(const std::set<int> &A) printf("\n"); } +void function_info::find_dominators( ) +{ + // find dominators using algorithm of Muchnick's Adv. Compiler Design & Implemmntation Fig 7.14 + printf("GPGPU-Sim PTX: Finding dominators for \'%s\'...\n", m_name.c_str() ); + fflush(stdout); + assert( m_basic_blocks.size() >= 2 ); // must have a distinquished entry block + std::vector<basic_block_t*>::iterator bb_itr = m_basic_blocks.begin(); + (*bb_itr)->dominator_ids.insert((*bb_itr)->bb_id); // the only dominator of the entry block is the entry + //copy all basic blocks to all dominator lists EXCEPT for the entry block + for (++bb_itr;bb_itr != m_basic_blocks.end(); bb_itr++) { + for (unsigned i = 0; i < m_basic_blocks.size(); i++) + (*bb_itr)->dominator_ids.insert(i); + } + bool change = true; + while (change) { + change = false; + for ( int h = 1/*skip entry*/; h < m_basic_blocks.size(); ++h ) { + assert( m_basic_blocks[h]->bb_id == (unsigned)h ); + std::set<int> T; + for (unsigned i=0;i< m_basic_blocks.size();i++) + T.insert(i); + for ( std::set<int>::iterator s = m_basic_blocks[h]->predecessor_ids.begin();s != m_basic_blocks[h]->predecessor_ids.end();s++) + intersect(T, m_basic_blocks[*s]->dominator_ids); + T.insert(h); + if (!is_equal(T, m_basic_blocks[h]->dominator_ids)) { + change = true; + m_basic_blocks[h]->dominator_ids = T; + } + } + } + //clean the basic block of dominators of it has no predecessors -- except for entry block + bb_itr = m_basic_blocks.begin(); + for (++bb_itr;bb_itr != m_basic_blocks.end(); bb_itr++) { + if ((*bb_itr)->predecessor_ids.empty()) + (*bb_itr)->dominator_ids.clear(); + } +} + void function_info::find_postdominators( ) { // find postdominators using algorithm of Muchnick's Adv. Compiler Design & Implemmntation Fig 7.14 @@ -567,6 +651,64 @@ void function_info::find_ipostdominators( ) // the exit node does not have an immediate post dominator, but everyone else should } +void function_info::find_idominators( ) +{ + // find immediate postdominator blocks, using algorithm of + // Muchnick's Adv. Compiler Design & Implemmntation Fig 7.15 + printf("GPGPU-Sim PTX: Finding immediate postdominators for \'%s\'...\n", m_name.c_str() ); + fflush(stdout); + assert( m_basic_blocks.size() >= 2 ); // must have a distinquished exit block + for (unsigned i=0; i<m_basic_blocks.size(); i++) { //initialize Tmp(n) to all pdoms of n except for n + m_basic_blocks[i]->Tmp_ids = m_basic_blocks[i]->dominator_ids; + assert( m_basic_blocks[i]->bb_id == i ); + m_basic_blocks[i]->Tmp_ids.erase(i); + } + for ( int n = 0; n < m_basic_blocks.size(); ++n) { + // point iterator to basic block before the exit + for( std::set<int>::iterator s=m_basic_blocks[n]->Tmp_ids.begin(); s != m_basic_blocks[n]->Tmp_ids.end(); s++ ) { + int bb_s = *s; + for( std::set<int>::iterator t=m_basic_blocks[n]->Tmp_ids.begin(); t != m_basic_blocks[n]->Tmp_ids.end(); ) { + std::set<int>::iterator t_next = t; t_next++; // might erase thing pointed to be t, invalidating iterator t + if( *s == *t ) { + t = t_next; + continue; + } + int bb_t = *t; + if( m_basic_blocks[bb_s]->dominator_ids.find(bb_t) != m_basic_blocks[bb_s]->dominator_ids.end() ) + m_basic_blocks[n]->Tmp_ids.erase(bb_t); + t = t_next; + } + } + } + unsigned num_idoms=0; + unsigned num_nopred = 0; + for ( int n = 0; n < m_basic_blocks.size(); ++n) { + assert( m_basic_blocks[n]->Tmp_ids.size() <= 1 ); + // if the above assert fails we have an error in either postdominator + // computation, the flow graph does not have a unique exit, or some other error + if( !m_basic_blocks[n]->Tmp_ids.empty() ) { + m_basic_blocks[n]->immediatedominator_id = *m_basic_blocks[n]->Tmp_ids.begin(); + num_idoms++; + } else if (m_basic_blocks[n]->predecessor_ids.empty()) { + num_nopred += 1; + } + } + assert( num_idoms == m_basic_blocks.size()-num_nopred ); + // the entry node does not have an immediate dominator, but everyone else should +} + +void function_info::print_dominators() +{ + printf("Printing dominators for function \'%s\':\n", m_name.c_str() ); + std::vector<int>::iterator bb_itr; + for (unsigned i = 0; i < m_basic_blocks.size(); i++) { + printf("ID: %d\t:", i); + for( std::set<int>::iterator j=m_basic_blocks[i]->dominator_ids.begin(); j!=m_basic_blocks[i]->dominator_ids.end(); j++) + printf(" %d", *j ); + printf("\n"); + } +} + void function_info::print_postdominators() { printf("Printing postdominators for function \'%s\':\n", m_name.c_str() ); @@ -589,6 +731,16 @@ void function_info::print_ipostdominators() } } +void function_info::print_idominators() +{ + printf("Printing immediate dominators for function \'%s\':\n", m_name.c_str() ); + std::vector<int>::iterator bb_itr; + for (unsigned i = 0; i < m_basic_blocks.size(); i++) { + printf("ID: %d\t:", i); + printf("%d\n", m_basic_blocks[i]->immediatedominator_id); + } +} + unsigned function_info::get_num_reconvergence_pairs() { if (!num_reconvergence_pairs) { @@ -703,7 +855,7 @@ unsigned type_info_key::type_decode( int type, size_t &size, int &basic_type ) arg_buffer_t copy_arg_to_buffer(ptx_thread_info * thread, operand_info actual_param_op, const symbol * formal_param) { if( actual_param_op.is_reg() ) { - ptx_reg_t value = thread->get_operand_value(actual_param_op); + ptx_reg_t value = thread->get_reg(actual_param_op.get_symbol()); return arg_buffer_t(formal_param,actual_param_op,value); } else if ( actual_param_op.is_param_local() ) { unsigned size=formal_param->get_size_in_bytes(); @@ -738,7 +890,7 @@ void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a) if( a.is_reg() ) { ptx_reg_t value = a.get_reg(); operand_info dst_reg = operand_info(a.get_dst()); - thread->set_operand_value(dst_reg,value); + thread->set_reg(dst_reg.get_symbol(),value); } else { const void *buffer = a.get_param_buffer(); size_t size = a.get_param_buffer_size(); @@ -807,6 +959,7 @@ static std::list<operand_info> check_operands( int opcode, ptx_instruction::ptx_instruction( int opcode, const symbol *pred, int neg_pred, + int pred_mod, symbol *label, const std::list<operand_info> &operands, const operand_info &return_var, @@ -822,6 +975,7 @@ ptx_instruction::ptx_instruction( int opcode, m_opcode = opcode; m_pred = pred; m_neg_pred = neg_pred; + m_pred_mod = pred_mod; m_label = label; const std::list<operand_info> checked_operands = check_operands(opcode,scalar_type,operands); m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() ); @@ -831,6 +985,9 @@ ptx_instruction::ptx_instruction( int opcode, m_hi = false; m_lo = false; m_uni = false; + m_exit = false; + m_abs = false; + m_neg = false; m_to_option = false; m_cache_option = 0; m_rounding_mode = RN_OPTION; @@ -941,6 +1098,15 @@ ptx_instruction::ptx_instruction( int opcode, break; case FTZ_OPTION: break; + case EXIT_OPTION: + m_exit = true; + break; + case ABS_OPTION: + m_abs = true; + break; + case NEG_OPTION: + m_neg = true; + break; case TO_OPTION: m_to_option = true; break; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 0887df0..43ac92a 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -283,6 +283,7 @@ public: return m_arch_reg_num; } void print_info(FILE *fp) const; + unsigned uid() const { return m_uid; } private: unsigned get_uid(); @@ -370,11 +371,21 @@ class operand_info { public: operand_info() { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = false; } operand_info( const symbol *addr ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; if ( addr->is_label() ) { @@ -402,8 +413,33 @@ public: m_neg_pred = false; m_is_return_var = false; } + operand_info( const symbol *addr1, const symbol *addr2 ) + { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; + m_uid = get_uid(); + m_valid = true; + m_type = memory_t; + m_value.m_vector_symbolic = new const symbol*[4]; + m_value.m_vector_symbolic[0] = addr1; + m_value.m_vector_symbolic[1] = addr2; + m_value.m_vector_symbolic[2] = NULL; + m_value.m_vector_symbolic[3] = NULL; + m_addr_offset = 0; + m_vector = false; + m_neg_pred = false; + m_is_return_var = false; + } operand_info( int builtin_id, int dim_mod ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -415,6 +451,11 @@ public: } operand_info( const symbol *addr, int offset ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -426,6 +467,11 @@ public: } operand_info( unsigned x ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -437,6 +483,11 @@ public: } operand_info( int x ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -448,6 +499,11 @@ public: } operand_info( float x ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -459,6 +515,11 @@ public: } operand_info( double x ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = false; @@ -470,6 +531,11 @@ public: } operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ) { + m_addr_space = 0; + m_operand_lohi = 0; + m_double_operand_type = 0; + m_operand_neg = false; + m_const_mem_offset = 0; m_uid = get_uid(); m_valid = true; m_vector = true; @@ -506,7 +572,9 @@ public: const symbol* vec_symbol(int idx) const { assert(idx < 4); - return m_value.m_vector_symbolic[idx]; + const symbol *result = m_value.m_vector_symbolic[idx]; + assert( result != NULL ); + return result; } const std::string &vec_name1() const @@ -564,12 +632,13 @@ public: int arch_reg_num(unsigned n) const { return (m_value.m_vector_symbolic[n])? m_value.m_vector_symbolic[n]->arch_reg_num() : -1; } bool is_label() const { return m_type == label_t;} bool is_builtin() const { return m_type == builtin_t;} + + // Memory operand used in ld / st instructions (ex. [__var1]) bool is_memory_operand() const { return m_type == memory_t;} + // Memory operand with immediate access (ex. s[0x0004] or g[$r1+=0x0004]) bool is_memory_operand2() const { - /* TODO: modify after integrate with ptxplus*/ - return m_type == memory_t; - /*return (m_addr_space==1 || m_addr_space==2 || m_addr_space==3 || m_addr_space==4);*/ + return (m_addr_space==1 || m_addr_space==2 || m_addr_space==3 || m_addr_space==4); } bool is_literal() const { return m_type == int_t || @@ -628,7 +697,16 @@ public: bool is_neg_pred() const { return m_neg_pred; } bool is_valid() const { return m_valid; } - int get_double_operand_type() const { return 0; /* TODO: modify after integrate with ptxplus*/ } + void set_addr_space(int set_value) { m_addr_space = set_value; } + int get_addr_space() const { return m_addr_space; } + void set_operand_lohi(int set_value) { m_operand_lohi = set_value; } + int get_operand_lohi() const { return m_operand_lohi; } + void set_double_operand_type(int set_value) { m_double_operand_type = set_value; } + int get_double_operand_type() const { return m_double_operand_type; } + void set_operand_neg() { m_operand_neg = true; } + bool get_operand_neg() const { return m_operand_neg; } + void set_const_mem_offset(addr_t set_value) { m_const_mem_offset = set_value; } + addr_t get_const_mem_offset() const { return m_const_mem_offset; } private: unsigned m_uid; @@ -636,6 +714,11 @@ private: bool m_vector; enum operand_type m_type; + int m_addr_space; + int m_operand_lohi; + int m_double_operand_type; + bool m_operand_neg; + addr_t m_const_mem_offset; union { int m_int; unsigned int m_unsigned; @@ -645,7 +728,7 @@ private: unsigned int m_vunsigned[4]; float m_vfloat[4]; double m_vdouble[4]; - const symbol* m_symbolic; + const symbol* m_symbolic; const symbol** m_vector_symbolic; } m_value; @@ -669,6 +752,7 @@ struct basic_block_t { is_entry=entry; is_exit=ex; immediatepostdominator_id = -1; + immediatedominator_id = -1; } ptx_instruction* ptx_begin; @@ -679,9 +763,20 @@ struct basic_block_t { std::set<int> dominator_ids; std::set<int> Tmp_ids; int immediatepostdominator_id; + int immediatedominator_id; bool is_entry; bool is_exit; unsigned bb_id; + + // if this basic block dom B + bool dom(const basic_block_t *B) { + return (B->dominator_ids.find(this->bb_id) != B->dominator_ids.end()); + } + + // if this basic block pdom B + bool pdom(const basic_block_t *B) { + return (B->postdominator_ids.find(this->bb_id) != B->postdominator_ids.end()); + } }; struct gpgpu_recon_t { @@ -694,6 +789,7 @@ public: ptx_instruction( int opcode, const symbol *pred, int neg_pred, + int pred_mod, symbol *label, const std::list<operand_info> &operands, const operand_info &return_var, @@ -706,7 +802,7 @@ public: void print_insn() const; void print_insn( FILE *fp ) const; unsigned uid() const { return m_uid;} - int get_opcode() { return m_opcode;} + int get_opcode() const { return m_opcode;} const char *get_opcode_cstr() const { if ( m_opcode != -1 ) { @@ -721,6 +817,7 @@ public: bool has_pred() const { return m_pred != NULL;} operand_info get_pred() const { return operand_info( m_pred );} bool get_pred_neg() const { return m_neg_pred;} + int get_pred_mod() const { return m_pred_mod;} const char *get_source() const { return m_source.c_str();} typedef std::vector<operand_info>::const_iterator const_iterator; @@ -827,6 +924,9 @@ public: bool is_lo() const { return m_lo;} bool is_wide() const { return m_wide;} bool is_uni() const { return m_uni;} + bool is_exit() const { return m_exit;} + bool is_abs() const { return m_abs;} + bool is_neg() const { return m_neg;} bool is_to() const { return m_to_option; } unsigned cache_option() const { return m_cache_option; } unsigned rounding_mode() const { return m_rounding_mode;} @@ -841,10 +941,22 @@ public: bool has_memory_read() const { if( m_opcode == LD_OP || m_opcode == TEX_OP ) return true; + // Source operands are memory operands + ptx_instruction::const_iterator op=op_iter_begin(); + for ( int n=0; op != op_iter_end(); op++, n++ ) { //process operands + if( n > 0 && op->is_memory_operand2()) // source operands only + return true; + } return false; } bool has_memory_write() const { if( m_opcode == ST_OP ) return true; + // Destination operand is a memory operand + ptx_instruction::const_iterator op=op_iter_begin(); + for ( int n=0; (op!=op_iter_end() && n<1); op++, n++ ) { //process operands + if( n==0 && op->is_memory_operand2()) // source operands only + return true; + } return false; } @@ -859,6 +971,7 @@ private: const symbol *m_pred; bool m_neg_pred; + int m_pred_mod; int m_opcode; const symbol *m_label; std::vector<operand_info> m_operands; @@ -868,7 +981,10 @@ private: bool m_wide; bool m_hi; bool m_lo; - bool m_uni; //if branch instruction, this evaluates to true for uniform branches (ie jumps) + bool m_exit; + bool m_abs; + bool m_neg; + bool m_uni; //if branch instruction, this evaluates to true for uniform branches (ie jumps) bool m_to_option; unsigned m_cache_option; unsigned m_rounding_mode; @@ -900,6 +1016,8 @@ public: m_value_set = true; m_value = v; } + void add_offset( unsigned offset ) { m_offset = offset; } + unsigned get_offset() { assert(m_valid); return m_offset; } std::string get_name() const { assert(m_valid); return m_name; } int get_type() const { assert(m_valid); return m_type; } param_t get_value() const { assert(m_value_set); return m_value; } @@ -911,6 +1029,7 @@ private: size_t m_size; bool m_value_set; param_t m_value; + unsigned m_offset; }; class function_info { @@ -943,22 +1062,31 @@ public: void print_basic_block_links(); void print_basic_block_dot(); + operand_info* find_break_target( ptx_instruction * p_break_insn ); //find the target of a break instruction void connect_basic_blocks( ); //iterate across m_basic_blocks of function, connecting basic blocks together + bool connect_break_targets(); //connecting break instructions with proper targets + + //iterate across m_basic_blocks of function, + //finding dominator blocks, using algorithm of + //Muchnick's Adv. Compiler Design & Implemmntation Fig 7.14 + void find_dominators( ); + void print_dominators(); + void find_idominators(); + void print_idominators(); //iterate across m_basic_blocks of function, //finding postdominator blocks, using algorithm of //Muchnick's Adv. Compiler Design & Implemmntation Fig 7.14 void find_postdominators( ); + void print_postdominators(); //iterate across m_basic_blocks of function, //finding immediate postdominator blocks, using algorithm of //Muchnick's Adv. Compiler Design & Implemmntation Fig 7.15 void find_ipostdominators( ); - - void print_postdominators(); - void print_ipostdominators(); + unsigned get_num_reconvergence_pairs(); void get_reconvergence_pairs(gpgpu_recon_t* recon_points); @@ -1032,6 +1160,7 @@ public: } void finalize( memory_space *param_mem ); + void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; const struct gpgpu_ptx_sim_kernel_info* get_kernel_info () { @@ -1292,9 +1421,10 @@ extern std::map<std::string,symbol_table*> g_sym_name_to_symbol_table; #define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) #define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM) #define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM) - #define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) + + extern bool g_keep_intermediate_files; void gpgpu_ptx_assemble( std::string kname, void *kinfo ); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 05ce397..d4a892f 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -94,6 +94,7 @@ type_info *g_var_type = NULL; // instruction definition stuff: const symbol *g_pred; int g_neg_pred; +int g_pred_mod; symbol *g_label; int g_opcode = -1; std::list<operand_info> g_operands; @@ -165,6 +166,7 @@ void init_instruction_state() DPRINTF("init_instruction_state"); g_pred = NULL; g_neg_pred = 0; + g_pred_mod = -1; g_label = NULL; g_opcode = -1; g_options.clear(); @@ -291,7 +293,8 @@ void add_instruction() DPRINTF("add_instruction: %s", ((g_opcode>0)?g_opcode_string[g_opcode]:"<label>") ); ptx_instruction *i = new ptx_instruction( g_opcode, g_pred, - g_neg_pred, + g_neg_pred, + g_pred_mod, g_label, g_operands, g_return_var, @@ -495,6 +498,21 @@ void add_identifier( const char *identifier, int array_dim, unsigned array_ident } } +void add_constptr(const char* identifier1, const char* identifier2, int offset) +{ + symbol *s1 = g_current_symbol_table->lookup(identifier1); + const symbol *s2 = g_current_symbol_table->lookup(identifier2); + parse_assert( s1 != NULL, "'from' constant identifier does not exist."); + parse_assert( s1 != NULL, "'to' constant identifier does not exist."); + + unsigned addr = s2->get_address(); + + printf("GPGPU-Sim PTX: moving \"%s\" from 0x%x to 0x%x (%s+%x)\n", + identifier1, s1->get_address(), addr+offset, identifier2, offset); + + s1->set_address( addr + offset ); +} + void add_function_arg() { if( g_func_info ) { @@ -570,7 +588,7 @@ void add_opcode( int opcode ) g_opcode = opcode; } -void add_pred( const char *identifier, int neg ) +void add_pred( const char *identifier, int neg, int predModifier ) { DPRINTF("add_pred"); const symbol *s = g_current_symbol_table->lookup(identifier); @@ -580,6 +598,7 @@ void add_pred( const char *identifier, int neg ) } g_pred = s; g_neg_pred = neg; + g_pred_mod = predModifier; } void add_option( int option ) @@ -588,6 +607,19 @@ void add_option( int option ) g_options.push_back( option ); } +void add_double_operand( const char *d1, const char *d2 ) +{ + //operands that access two variables. + //eg. s[$ofs1+$r0], g[$ofs1+=$r0] + //TODO: Not sure if I'm going to use this for storing to two destinations or not. + + DPRINTF("add_double_operand"); + const symbol *s1 = g_current_symbol_table->lookup(d1); + const symbol *s2 = g_current_symbol_table->lookup(d2); + parse_assert( s1 != NULL && s2 != NULL, "component(s) missing declarations."); + g_operands.push_back( operand_info(s1,s2) ); +} + void add_2vector_operand( const char *d1, const char *d2 ) { DPRINTF("add_2vector_operand"); @@ -631,6 +663,109 @@ void add_memory_operand() g_operands.back().make_memory_operand(); } +/*TODO: add other memory locations*/ +void change_memory_addr_space(const char *identifier) +{ + /*0 = N/A, not reading from memory + *1 = global memory + *2 = shared memory + *3 = const memory segment + *4 = local memory segment + */ + + bool recognizedType = false; + + DPRINTF("change_memory_addr_space"); + assert( !g_operands.empty() ); + if(!strcmp(identifier, "g")) + { + g_operands.back().set_addr_space(1); + recognizedType = true; + } + if(!strcmp(identifier, "s")) + { + g_operands.back().set_addr_space(2); + recognizedType = true; + } + // For constants, check if the first character is 'c' + char c[2]; + strncpy(c, identifier, 1); c[1] = '\0'; + if(!strcmp(c, "c")) + { + g_operands.back().set_addr_space(3); + parse_assert(g_current_symbol_table->lookup(identifier) != NULL, "Constant was not defined."); + g_operands.back().set_const_mem_offset(g_current_symbol_table->lookup(identifier)->get_address()); + recognizedType = true; + } + // For local memory, check if the first character is 'l' + char l[2]; + strncpy(l, identifier, 1); l[1] = '\0'; + if(!strcmp(l, "l")) + { + g_operands.back().set_addr_space(4); + parse_assert(g_current_symbol_table->lookup(identifier) != NULL, "Local memory segment was not defined."); + g_operands.back().set_const_mem_offset(g_current_symbol_table->lookup(identifier)->get_address()); + recognizedType = true; + } + + parse_assert(recognizedType, "Error: unrecognized memory type."); +} + +void change_operand_lohi( int lohi ) +{ + /*0 = N/A, read entire operand + *1 = lo, reading from lowest bits + *2 = hi, reading from highest bits + */ + + DPRINTF("change_operand_lohi"); + assert( !g_operands.empty() ); + + g_operands.back().set_operand_lohi(lohi); + +} + +void change_double_operand_type( int operand_type ) +{ + /* + *-3 = reg / reg (set instruction, but both get same value) + *-2 = reg | reg (cvt instruction) + *-1 = reg | reg (set instruction) + *0 = N/A, default + *1 = reg + reg + *2 = reg += reg + *3 = reg += immediate + */ + + DPRINTF("change_double_operand_type"); + assert( !g_operands.empty() ); + + // For double destination operands, ensure valid instruction + if( operand_type == -1 || operand_type == -2 ) { + if((g_opcode == SET_OP)||(g_opcode == SETP_OP)) + g_operands.back().set_double_operand_type(-1); + else + g_operands.back().set_double_operand_type(-2); + } else if( operand_type == -3 ) { + if(g_opcode == SET_OP) + g_operands.back().set_double_operand_type(operand_type); + else + parse_assert(0, "Error: Unsupported use of double destination operand."); + } else { + g_operands.back().set_double_operand_type(operand_type); + } + +} + +void change_operand_neg( ) +{ + DPRINTF("change_operand_neg"); + assert( !g_operands.empty() ); + + g_operands.back().set_operand_neg(); + +} + void add_literal_int( int value ) { DPRINTF("add_literal_int"); @@ -741,3 +876,11 @@ void add_pragma( const char *str ) { printf("GPGPU-Sim PTX: Warning -- ignoring pragma '%s'\n", str ); } + +void version_header(double a) {} //intentional dummy function +void target_header(char* a) {} //intentional dummy function +void target_header2(char* a, char* b) {} //intentional dummy function + +void func_header(char* a) {} //intentional dummy function +void func_header_info(char* a) {} //intentional dummy function +void func_header_info_int(char* a, int b) {} //intentional dummy function diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 8ffe593..4852e39 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -90,7 +90,7 @@ void add_neg_pred_operand( const char *identifier ); void add_variables(); void set_variable_type(); void add_opcode( int opcode ); -void add_pred( const char *identifier, int negate ); +void add_pred( const char *identifier, int negate, int predModifier ); void add_2vector_operand( const char *d1, const char *d2 ); void add_3vector_operand( const char *d1, const char *d2, const char *d3 ); void add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 ); @@ -114,6 +114,18 @@ void add_version_info( float ver ); void *reset_symtab(); void set_symtab(void*); void add_pragma( const char *str ); +void func_header(char* a); +void func_header_info(char* a); +void func_header_info_int(char* a, int b); +void add_constptr(const char* identifier1, const char* identifier2, int offset); +void target_header(char* a); +void target_header2(char* a, char* b); +void add_double_operand( const char *d1, const char *d2 ); +void change_memory_addr_space( const char *identifier ); +void change_operand_lohi( int lohi ); +void change_double_operand_type( int addr_type ); +void change_operand_neg( ); +void version_header(double a); #ifdef __cplusplus } #endif diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 685aedb..6dea2e3 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -238,6 +238,8 @@ ptx_thread_info::ptx_thread_info() m_last_dram_callback.function = NULL; m_last_dram_callback.instruction = NULL; m_regs.push_back( reg_map_t() ); + m_debug_trace_regs_modified.push_back( reg_map_t() ); + m_debug_trace_regs_read.push_back( reg_map_t() ); m_callstack.push_back( stack_entry() ); m_RPC = -1; m_RPC_updated = false; @@ -266,7 +268,12 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod ) return (unsigned)(gpu_sim_cycle + gpu_tot_sim_cycle); case CLOCK64_REG: abort(); // change return value to unsigned long long? - return gpu_sim_cycle + gpu_tot_sim_cycle; + // GPGPUSim clock is 4 times slower - multiply by 4 + return (gpu_sim_cycle + gpu_tot_sim_cycle)*4; + case HALFCLOCK_ID: + // GPGPUSim clock is 4 times slower - multiply by 4 + // Hardware clock counter is incremented at half the shader clock frequency - divide by 2 (Henry '10) + return (gpu_sim_cycle + gpu_tot_sim_cycle)*2; case CTAID_REG: assert( dim_mod < 3 ); return m_ctaid[dim_mod]; @@ -305,6 +312,18 @@ void ptx_thread_info::set_info( function_info *func ) m_PC = func->get_start_PC(); } +void ptx_thread_info::cpy_tid_to_reg( int x, int y, int z) +{ + //copies %tid.x, %tid.y and %tid.z into $r0 + ptx_reg_t data; + data.s64=0; + + data.u32=(x + (y<<16) + (z<<26)); + + const symbol *r0 = m_symbol_table->lookup("$r0"); + set_reg(r0,data); +} + void ptx_thread_info::print_insn( unsigned pc, FILE * fp ) const { m_func_info->print_insn(pc,fp); @@ -346,6 +365,7 @@ static void print_reg( FILE *fp, std::string name, ptx_reg_t value, symbol_table fprintf( fp, "non-scalar type\n" ); break; } + fflush(fp); } static void print_reg( std::string name, ptx_reg_t value, symbol_table *symtab ) @@ -361,6 +381,8 @@ void ptx_thread_info::callstack_push( unsigned pc, unsigned rpc, const symbol *r assert( m_func_info != NULL ); m_callstack.push_back( stack_entry(m_symbol_table,m_func_info,pc,rpc,return_var_src,return_var_dst,call_uid) ); m_regs.push_back( reg_map_t() ); + m_debug_trace_regs_modified.push_back( reg_map_t() ); + m_debug_trace_regs_read.push_back( reg_map_t() ); m_local_mem_stack_pointer += m_func_info->local_mem_framesize(); } @@ -387,6 +409,8 @@ bool ptx_thread_info::callstack_pop() } m_callstack.pop_back(); m_regs.pop_back(); + m_debug_trace_regs_modified.pop_back(); + m_debug_trace_regs_read.pop_back(); // write return value into caller frame if( rv_dst != NULL ) @@ -441,38 +465,63 @@ const ptx_instruction *ptx_thread_info::get_inst( addr_t pc ) const return m_func_info->get_instruction(pc); } -void ptx_thread_info::dump_regs() -{ - dump_regs(stdout); -} - void ptx_thread_info::dump_regs( FILE *fp ) { + if(m_regs.empty()) return; + if(m_regs.back().empty()) return; fprintf(fp,"Register File Contents:\n"); + fflush(fp); reg_map_t::const_iterator r; for ( r=m_regs.back().begin(); r != m_regs.back().end(); ++r ) { - std::string name = r->first->name(); + const symbol *sym = r->first; ptx_reg_t value = r->second; + std::string name = sym->name(); print_reg(name,value,m_symbol_table); } } -void ptx_thread_info::dump_modifiedregs() +void ptx_thread_info::dump_modifiedregs(FILE *fp) { - dump_modifiedregs(stdout); + if( !(m_debug_trace_regs_modified.empty() || + m_debug_trace_regs_modified.back().empty()) ) { + fprintf(fp,"Output Registers:\n"); + fflush(fp); + reg_map_t::iterator r; + for ( r=m_debug_trace_regs_modified.back().begin(); r != m_debug_trace_regs_modified.back().end(); ++r ) { + const symbol *sym = r->first; + std::string name = sym->name(); + ptx_reg_t value = r->second; + print_reg(fp,name,value,m_symbol_table); + } + } + if( !(m_debug_trace_regs_read.empty() || + m_debug_trace_regs_read.back().empty()) ) { + fprintf(fp,"Input Registers:\n"); + fflush(fp); + reg_map_t::iterator r; + for ( r=m_debug_trace_regs_read.back().begin(); r != m_debug_trace_regs_read.back().end(); ++r ) { + const symbol *sym = r->first; + std::string name = sym->name(); + ptx_reg_t value = r->second; + print_reg(fp,name,value,m_symbol_table); + } + } } -void ptx_thread_info::dump_modifiedregs(FILE *fp) +void ptx_thread_info::push_breakaddr(const operand_info &breakaddr) { - if( m_debug_trace_regs_modified.empty() ) - return; - fprintf(fp,"Modified Registers:\n"); - reg_map_t::const_iterator r; - for ( r=m_debug_trace_regs_modified.begin(); r != m_debug_trace_regs_modified.end(); ++r ) { - std::string name = r->first->name(); - ptx_reg_t value = r->second; - print_reg(name,value,m_symbol_table); + m_breakaddrs.push(breakaddr); +} + +const operand_info& ptx_thread_info::pop_breakaddr() +{ + if(m_breakaddrs.empty()) { + printf("empty breakaddrs stack"); + assert(0); } + operand_info& breakaddr = m_breakaddrs.top(); + m_breakaddrs.pop(); + return breakaddr; } void ptx_thread_info::set_npc( const function_info *f ) @@ -482,6 +531,7 @@ void ptx_thread_info::set_npc( const function_info *f ) m_symbol_table = m_func_info->get_symtab(); } + void feature_not_implemented( const char *f ) { printf("GPGPU-Sim: feature '%s' not supported\n", f ); diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index 8730f2b..4741769 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -105,6 +105,10 @@ struct param_t { size_t offset; }; +#include <stack> + +#include "memory.h" + union ptx_reg_t { ptx_reg_t() { bits.ms = 0; @@ -155,7 +159,13 @@ union ptx_reg_t { unsigned ls; unsigned ms; } bits; - unsigned pred : 1; + struct { + unsigned int lowest; + unsigned int low; + unsigned int high; + unsigned int highest; + } u128; + unsigned pred : 4; }; @@ -238,7 +248,7 @@ public: m_ptx_extensions = extensions; } float ver() const { assert(m_valid); return m_ptx_version; } - float extensions() const { assert(m_valid); return m_ptx_extensions; } + unsigned extensions() const { assert(m_valid); return m_ptx_extensions; } private: bool m_valid; float m_ptx_version; @@ -251,10 +261,11 @@ public: ptx_thread_info(); const ptx_version &get_ptx_version() const; - ptx_reg_t get_operand_value( const symbol *reg ); - ptx_reg_t get_operand_value( const operand_info &op ); - void set_operand_value( const operand_info &dst, const ptx_reg_t &data ); - void set_operand_value( const symbol *dst, const ptx_reg_t &data ); + void set_reg( const symbol *reg, const ptx_reg_t &value ); + ptx_reg_t get_reg( const symbol *reg ); + ptx_reg_t get_operand_value( const operand_info &op, operand_info dstInfo, unsigned opType, ptx_thread_info *thread, int derefFlag ); + void set_operand_value( const operand_info &dst, const ptx_reg_t &data, unsigned type, ptx_thread_info *thread, const ptx_instruction *pI ); + void set_operand_value( const operand_info &dst, const ptx_reg_t &data, unsigned type, ptx_thread_info *thread, const ptx_instruction *pI, int overflow, int carry ); void get_vector_operand_values( const operand_info &op, ptx_reg_t* ptx_regs, unsigned num_elements ); void set_vector_operand_values( const operand_info &dst, const ptx_reg_t &data1, @@ -345,6 +356,7 @@ public: m_tid[1] = y; m_tid[2] = z; } + void cpy_tid_to_reg( int x, int y, int z); void set_ctaid( int x, int y, int z) { m_ctaid[0] = x; @@ -413,14 +425,13 @@ public: { m_PC = m_NPC; } - void dump_regs(); - void dump_regs(FILE *fp); - void dump_modifiedregs(); + void dump_regs(FILE * fp); void dump_modifiedregs(FILE *fp); - void clear_modifiedregs() { m_debug_trace_regs_modified.clear();} + void clear_modifiedregs() { m_debug_trace_regs_modified.back().clear(); m_debug_trace_regs_read.back().clear(); } function_info *get_finfo() { return m_func_info; } const function_info *get_finfo() const { return m_func_info; } - + void push_breakaddr(const operand_info &breakaddr); + const operand_info& pop_breakaddr(); void enable_debug_trace() { m_enable_debug_trace = true; } unsigned get_local_mem_stack_pointer() const { return m_local_mem_stack_pointer; } @@ -435,6 +446,8 @@ public: ptx_reg_t m_last_set_operand_value; private: + + unsigned m_uid; core_t *m_core; bool m_valid; @@ -468,9 +481,11 @@ private: typedef tr1_hash_map<const symbol*,ptx_reg_t> reg_map_t; std::list<reg_map_t> m_regs; - + std::list<reg_map_t> m_debug_trace_regs_modified; + std::list<reg_map_t> m_debug_trace_regs_read; bool m_enable_debug_trace; - reg_map_t m_debug_trace_regs_modified; // track the modified register for each executed insn + + std::stack<class operand_info> m_breakaddrs; }; addr_t generic_to_local( unsigned smid, unsigned hwtid, addr_t addr ); diff --git a/src/debug.cc b/src/debug.cc index ac0daec..1f5388b 100644 --- a/src/debug.cc +++ b/src/debug.cc @@ -123,7 +123,7 @@ void gpgpu_debug() } else if( !strcmp(tok,"q") || !strcmp(tok,"quit") ) { printf("\nreally quit GPGPU-Sim (y/n)?\n"); fgets(line,1024,stdin); - tok = strtok(NULL," \t\n"); + tok = strtok(line," \t\n"); if( !strcmp(tok,"y") ) { exit(0); } else { diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc index 4ebcee9..278ada4 100644 --- a/src/gpgpu-sim/scoreboard.cc +++ b/src/gpgpu-sim/scoreboard.cc @@ -34,8 +34,6 @@ void Scoreboard::printContents() { // Mark register as write-pending void Scoreboard::reserveRegister(int wid, int regnum) { - //printf("Reserving register sid=%d wid=%d regnum=%d\n", sid, wid, regnum); - if( !(reg_table[wid].find(regnum) == reg_table[wid].end()) ){ printf("Error: trying to reserve an already reserved register (sid=%d, wid=%d, regnum=%d).", sid, wid, regnum); assert(reg_table[wid].find(regnum) == reg_table[wid].end()); @@ -47,8 +45,6 @@ void Scoreboard::reserveRegister(int wid, int regnum) { // Unmark register as write-pending void Scoreboard::releaseRegister(int wid, int regnum) { - //printf("Releasing register sid=%d wid=%d regnum=%d\n", sid, wid, regnum); - if( !(reg_table[wid].find(regnum) != reg_table[wid].end()) ) { printf("Error: trying to release an unreserved register (sid=%d, wid=%d, regnum=%d).", sid, wid, regnum); assert(reg_table[wid].find(regnum) != reg_table[wid].end()); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 2f304b7..5cd7d7f 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -92,7 +92,12 @@ void gpgpu_ptx_sim_init_perf() ptx_reg_options(opp); option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options - srand(1); + srand(1); + + // Open instructions debug output file for writing + if(g_ptx_inst_debug_to_file != 0) { + ptx_inst_debug_file = fopen(g_ptx_inst_debug_file, "w"); + } fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); option_parser_print(opp, stdout); |