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]

16 files changed, 1864 insertions, 463 deletions

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