From bb820c116764d7a1b8e071137d32b74e7f34dd2f Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Mon, 30 Aug 2010 11:05:15 -0800 Subject: integrating CL 7421-7434 from fermi-test this integrates the ptxplus functional model changes, but has only been verified with CUDA 3.1 PTX regressions (which pass) still need to verify ptxplus itself is working in this branch [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7435] --- src/cuda-sim/Makefile | 156 +-- src/cuda-sim/cuda-sim.cc | 140 ++- src/cuda-sim/cuda-sim.h | 1 + .../decuda_pred_table/decuda_pred_table.cc | 79 ++ src/cuda-sim/decuda_pred_table/decuda_pred_table.h | 1 + src/cuda-sim/instructions.cc | 1191 ++++++++++++++++---- src/cuda-sim/opcodes.def | 8 + src/cuda-sim/opcodes.h | 1 + src/cuda-sim/ptx.l | 24 + src/cuda-sim/ptx.y | 120 +- src/cuda-sim/ptx_ir.cc | 170 ++- src/cuda-sim/ptx_ir.h | 154 ++- src/cuda-sim/ptx_parser.cc | 147 ++- src/cuda-sim/ptx_parser.h | 14 +- src/cuda-sim/ptx_sim.cc | 86 +- src/cuda-sim/ptx_sim.h | 41 +- 16 files changed, 1867 insertions(+), 466 deletions(-) create mode 100644 src/cuda-sim/decuda_pred_table/decuda_pred_table.cc create mode 100644 src/cuda-sim/decuda_pred_table/decuda_pred_table.h (limited to 'src/cuda-sim') 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 TextureToArrayMap; // texture bindings std::map 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(); @@ -647,6 +675,35 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread, *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()){ + *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. @@ -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::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::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 *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 TextureToInfoMap; extern std::map 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(); - - 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(); + 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(); + } + + } 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; + + 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 { + finalResult = 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); } - } 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 { - name = op.name().c_str(); - printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); - assert(0); } - return result; + 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 ) +{ + 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 operand_info &dst, const ptx_reg_t &data, unsigned type, ptx_thread_info *thread, const ptx_instruction *pI ) { - m_regs.back()[ dst.get_symbol() ] = data; - if (m_enable_debug_trace ) { - m_debug_trace_regs_modified[ dst.get_symbol() ] = 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); + } } - m_last_set_operand_value = data; -} -void ptx_thread_info::set_operand_value( const symbol *dst, const ptx_reg_t &data ) -{ - m_regs.back()[ dst ] = data; - if (m_enable_debug_trace ) { - m_debug_trace_regs_modified[ dst ] = data; + // 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; + } + + // 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); } - m_last_set_operand_value = data; + + } 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); + 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, 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 + ((asrc1(); 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::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; { @@ -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 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 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 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::iterator bb_itr; + std::vector::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 orig_successor_ids = p_bb->successor_ids; + + // erase the previous linkage with old successors + for(std::set::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 &A, const std::set &B ) { // return intersection of A and B in A @@ -493,6 +539,44 @@ void print_set(const std::set &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::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 T; + for (unsigned i=0;i< m_basic_blocks.size();i++) + T.insert(i); + for ( std::set::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; iTmp_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::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::iterator t=m_basic_blocks[n]->Tmp_ids.begin(); t != m_basic_blocks[n]->Tmp_ids.end(); ) { + std::set::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::iterator bb_itr; + for (unsigned i = 0; i < m_basic_blocks.size(); i++) { + printf("ID: %d\t:", i); + for( std::set::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::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 check_operands( int opcode, ptx_instruction::ptx_instruction( int opcode, const symbol *pred, int neg_pred, + int pred_mod, symbol *label, const std::list &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 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 dominator_ids; std::set 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 &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::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 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 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 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]:"