integrating recent changes from fermi-test into fermi

(i'll use "fermi" for more disruptive changes to the pipeline model such
as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...)

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805]

20 files changed, 1167 insertions, 730 deletions

diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h
index 5abf4f3..367563f 100644
--- a/src/cuda-sim/cuda-math.h
+++ b/src/cuda-sim/cuda-math.h
@@ -47,6 +47,7 @@ namespace cuda_math {
 #define __attribute__(a) // to remove warnings inside math_functions.h
 #undef INT_MAX
 
+#if CUDART_VERSION < 3000
 // DEVICE_BUILTIN
    struct int4 {
       int x, y, z, w;
@@ -70,18 +71,245 @@ namespace cuda_math {
 
 extern float rsqrtf(float); // CUDA 2.3 beta
 
-#if CUDART_VERSION < 3000
 #define CUDA_FLOAT_MATH_FUNCTIONS
 #include <device_types.h>
 #define __CUDA_INTERNAL_COMPILATION__
 #include <math_functions.h>
 #undef __CUDA_INTERNAL_COMPILATION__
 #undef __attribute__
+
+// float to integer conversion 
+int float2int(float a, enum cudaRoundMode mode)
+{
+   return __internal_float2uint(a, mode); 
+}
+
+// float to unsigned integer conversion 
+unsigned int float2uint(float a, enum cudaRoundMode mode)
+{
+   return __internal_float2uint(a, mode); 
+}
+
+float __ll2float_rz(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TOWARDZERO); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ll2float_ru(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_UPWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ll2float_rd(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_DOWNWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+
 #else
+
 #define CUDA_FLOAT_MATH_FUNCTIONS
-#include "/home/taamodt/nvcuda/2.3/cuda/include/device_types.h"
+#define __CUDACC__
+
+// implementing int to float intrinsics with different rounding modes 
+#include <device_types.h>
+#include <fenv.h>
+
+// 32-bit integer to float
+float __int2float_rn(int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TONEAREST); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __int2float_rz(int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TOWARDZERO); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __int2float_ru(int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_UPWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __int2float_rd(int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_DOWNWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+
+// 32-bit unsigned integer to float
+float __uint2float_rn(unsigned int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TONEAREST); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __uint2float_rz(unsigned int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TOWARDZERO); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __uint2float_ru(unsigned int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_UPWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __uint2float_rd(unsigned int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_DOWNWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+
+// 64-bit integer to float
+float __ll2float_rn(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TONEAREST); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ll2float_rz(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TOWARDZERO); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ll2float_ru(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_UPWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ll2float_rd(long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_DOWNWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+
+// 64-bit unsigned integer to float 
+float __ull2float_rn(unsigned long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TONEAREST); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ull2float_rz(unsigned long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_TOWARDZERO); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ull2float_ru(unsigned long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_UPWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+float __ull2float_rd(unsigned long long int a) {
+   int orig_rnd_mode = fegetround();
+   fesetround(FE_DOWNWARD); 
+   float b = a;
+   fesetround(orig_rnd_mode); 
+   return b;
+}
+
+// float to integer conversion 
+int float2int(float a, enum cudaRoundMode mode)
+{
+   int tmp;
+   switch (mode) {
+   case cuda_math::cudaRoundZero: tmp = truncf(a);     break;
+   case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break;
+   case cuda_math::cudaRoundMinInf: tmp = floorf(a);     break;
+   case cuda_math::cudaRoundPosInf: tmp = ceilf(a);      break;
+   default: abort();
+   }
+   return tmp; 
+}
+
+int __internal_float2int(float a, enum cudaRoundMode mode) 
+{
+   return float2int(a, mode); 
+}
+
+// float to unsigned integer conversion 
+unsigned int float2uint(float a, enum cudaRoundMode mode)
+{
+   unsigned int tmp;
+   switch (mode) {
+   case cuda_math::cudaRoundZero: tmp = truncf(a);     break;
+   case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break;
+   case cuda_math::cudaRoundMinInf: tmp = floorf(a);     break;
+   case cuda_math::cudaRoundPosInf: tmp = ceilf(a);      break;
+   default: abort();
+   }
+   return tmp; 
+}
+
+unsigned int __internal_float2uint(float a, enum cudaRoundMode mode) 
+{
+   return float2uint(a, mode); 
+}
+
+// intrinsic for division 
+float fdividef(float a, float b)
+{
+   return (a / b); 
+}
+
+float __internal_accurate_fdividef(float a, float b)
+{
+   return fdividef(a, b); 
+}
+
+// intrinsic for saturate  (clamp values beyond 0 and 1)
+float __saturatef(float a)
+{
+   float b; 
+   if (isnan(a)) b = 0.0f; 
+   else if (a >= 1.0f) b = 1.0f;
+   else if (a <= 0.0f) b = 0.0f; 
+   else b = a; 
+   return b; 
+}
+
+// intrinsic for power 
+float __powf(float a, float b)
+{
+   return powf(a, b);
+}
+
+#undef __CUDACC__
 #define __CUDA_INTERNAL_COMPILATION__
-#include "/home/taamodt/nvcuda/2.3/cuda/include/math_functions.h"
+#include <math_functions.h>
 #undef __CUDA_INTERNAL_COMPILATION__
 #undef __attribute__
 #endif
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 0a6ccd6..de4445a 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -99,7 +99,6 @@ std::map<const struct textureReference*,const struct cudaArray*> TextureToArrayM
 std::map<const struct textureReference*, const struct textureInfo*> TextureToInfoMap;
 std::map<std::string, const struct textureReference*> NameToTextureMap;
 unsigned int g_texcache_linesize;
-int gpgpu_option_spread_blocks_across_cores = 0;
 unsigned gpgpu_param_num_shaders = 0;
 
 void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref)
@@ -208,10 +207,12 @@ int gpgpu_ptx_sim_sizeofTexture(const char* name)
    return array->size;
 }
 
-unsigned g_assemble_code_next_pc=1; 
+unsigned g_assemble_code_next_pc=0; 
 std::map<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;
 
+#define MAX_INST_SIZE 8 /*bytes*/
+
 void function_info::ptx_assemble()
 {
    if( m_assembled ) {
@@ -220,16 +221,22 @@ void function_info::ptx_assemble()
 
    // get the instructions into instruction memory...
    unsigned num_inst = m_instructions.size();
-   m_instr_mem = new ptx_instruction*[ num_inst ];
-   m_instr_mem_size = num_inst;
+   m_instr_mem_size = MAX_INST_SIZE*(num_inst+1);
+   m_instr_mem = new ptx_instruction*[ m_instr_mem_size ];
 
    printf("GPGPU-Sim PTX: instruction assembly for function \'%s\'... ", m_name.c_str() );
    fflush(stdout);
    std::list<ptx_instruction*>::iterator i;
-   addr_t n=0; // offset in m_instr_mem
+
    addr_t PC = g_assemble_code_next_pc; // globally unique address (across functions)
+   // start function on an aligned address
+   for( unsigned i=0; i < (PC%MAX_INST_SIZE); i++ ) 
+      s_g_pc_to_insn.push_back((ptx_instruction*)NULL);
+   PC += PC%MAX_INST_SIZE; 
    m_start_PC = PC;
-   s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + m_instructions.size());
+
+   addr_t n=0; // offset in m_instr_mem
+   s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size());
    for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) {
       ptx_instruction *pI = *i;
       if ( pI->is_label() ) {
@@ -239,17 +246,22 @@ void function_info::ptx_assemble()
          g_pc_to_finfo[PC] = this;
          m_instr_mem[n] = pI;
          s_g_pc_to_insn.push_back(pI);
-         assert(pI == s_g_pc_to_insn[PC - 1]);
+         assert(pI == s_g_pc_to_insn[PC]);
          pI->set_m_instr_mem_index(n);
          pI->set_PC(PC);
-         n++;
-         PC++;
+         assert( pI->inst_size() <= MAX_INST_SIZE );
+         for( unsigned i=1; i < pI->inst_size(); i++ ) {
+            s_g_pc_to_insn.push_back((ptx_instruction*)NULL);
+            m_instr_mem[n+i]=NULL;
+         }
+         n  += pI->inst_size();
+         PC += pI->inst_size();
       }
    }
    g_assemble_code_next_pc=PC;
-   for ( unsigned ii=0; ii < n; ii++ ) { // handle branch instructions
+   for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions
       ptx_instruction *pI = m_instr_mem[ii];
-      if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP ) {
+      if ( pI->get_opcode() == BRA_OP || pI->get_opcode() == BREAKADDR_OP  || pI->get_opcode() == CALLP_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.",
@@ -262,6 +274,11 @@ void function_info::ptx_assemble()
          target.set_type(label_t);
       }
    }
+   for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions
+      ptx_instruction *pI = m_instr_mem[ii];
+      pI->pre_decode();
+   }
+
    printf("  done.\n");
    fflush(stdout);
 
@@ -328,12 +345,12 @@ bool isspace_shared( unsigned smid, addr_t addr )
 
 bool isspace_global( addr_t addr )
 {
-   return (addr > GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT);
+   return (addr >= GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT);
 }
 
 memory_space_t whichspace( addr_t addr )
 {
-   if( (addr > GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT) ) {
+   if( (addr >= GLOBAL_HEAP_START) || (addr < STATIC_ALLOC_LIMIT) ) {
       return global_space;
    } else if( addr > SHARED_GENERIC_START ) {
       return shared_space;
@@ -446,14 +463,6 @@ void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count )
    fflush(stdout);
 }
 
-int ptx_thread_done( void *thd )
-{
-   ptx_thread_info *the_thread = (ptx_thread_info *) thd;
-   int result = 0;
-   result = (the_thread==NULL) || the_thread->is_done();
-   return result;
-}
-
 const char * ptx_get_fname( unsigned PC )
 {
     static const char *null_ptr = "<null finfo ptr>";
@@ -471,14 +480,6 @@ unsigned ptx_thread_donecycle( void *thr )
    return the_thread->donecycle();
 }
 
-int ptx_thread_get_next_pc( void *thd )
-{
-   ptx_thread_info *the_thread = (ptx_thread_info *) thd;
-   if ( the_thread == NULL )
-      return -1;
-   return the_thread->get_pc(); // PC should already be updatd to next PC at this point (was set in shader_decode() last time thread ran)
-}
-
 void* ptx_thread_get_next_finfo( void *thd )
 {
    ptx_thread_info *the_thread = (ptx_thread_info *) thd;
@@ -531,7 +532,7 @@ void ptx_print_insn( address_type pc, FILE *fp )
 {
    std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
    if( f == g_pc_to_finfo.end() ) {
-       fprintf(fp,"<no instruction at address 0x%x (%u)>", pc, pc );
+       fprintf(fp,"<no instruction at address 0x%x>", pc );
        return;
    }
    function_info *finfo = f->second;
@@ -539,108 +540,118 @@ void ptx_print_insn( address_type pc, FILE *fp )
    finfo->print_insn(pc,fp);
 }
 
-static void get_opcode_info( const ptx_instruction *pI, unsigned opcode, unsigned *cycles, unsigned *op_type )
+static void get_opcode_info( const ptx_instruction *pI, unsigned opcode, unsigned *cycles, op_type *op )
 {
-   *op_type = ALU_OP;
+   *op = ALU_OP;
    *cycles = 1;
    if ( opcode == LD_OP ) {
-      *op_type = LOAD_OP;
+      *op = LOAD_OP;
    } else if ( opcode == ST_OP ) {
-      *op_type = STORE_OP;
+      *op = STORE_OP;
    } else if ( opcode == BRA_OP ) {
-      *op_type = BRANCH_OP;
+      *op = BRANCH_OP;
    } else if ( opcode == BREAKADDR_OP ) {
-      *op_type = BRANCH_OP;
+      *op = BRANCH_OP;
    } else if ( opcode == TEX_OP ) {
-      *op_type = LOAD_OP;
+      *op = LOAD_OP;
    } else if ( opcode == ATOM_OP ) {
-      *op_type = LOAD_OP; // make atomics behave more like a load.
+      *op = LOAD_OP; // timing model treats this like load for now
    } else if ( opcode == BAR_OP ) {
-      *op_type = BARRIER_OP;
-   }
+      *op = BARRIER_OP;
+   } else if ( opcode == MEMBAR_OP ) 
+      *op = MEMORY_BARRIER_OP;
 
    // Floating point instructions
-   if( opcode == RCP_OP || opcode == LG2_OP || opcode == RSQRT_OP ) {
-      *cycles = 4;
-      *op_type = SFU_OP;
+   if( opcode == RCP_OP ) {
+       *cycles = 2;
+       *op = SFU_OP;
+   } else if ( opcode == LG2_OP || opcode == RSQRT_OP ) {
+       *cycles = 4;
+       *op = SFU_OP;
    } else if( opcode == SQRT_OP || opcode == SIN_OP || opcode == COS_OP || opcode == EX2_OP ) {
       *cycles = 4;
-      *op_type = SFU_OP;
+      *op = SFU_OP;
    } else if( opcode == DIV_OP ) {
       // Floating point only
       if( pI->get_type() == F32_TYPE || pI->get_type() == F64_TYPE ) {
-         *cycles = 8;         
-         *op_type = SFU_OP;
+         *cycles = 4;         
+         *op = SFU_OP;
       }
    }
    // Integer instructions
    if( opcode == MUL_OP ) {
       if( pI->get_type() == B32_TYPE || pI->get_type() == U32_TYPE || pI->get_type() == S32_TYPE ) {
          // 32-bit integer instruction
-         *cycles = 4;
-         *op_type = SFU_OP;
+         *cycles = 5;
+         *op = SFU_OP;
       }
+      if( pI->get_type() == F32_TYPE || pI->get_type() == F64_TYPE ) 
+         *op = ALU_SFU_OP;
    }
+   if( opcode == MAD_OP ) {
+       if( pI->get_type() == B32_TYPE || pI->get_type() == U32_TYPE || pI->get_type() == S32_TYPE ) {
+          // 32-bit integer instruction
+          *cycles = 6;
+          *op = SFU_OP;
+       }
+    }
+}
+
+void ptx_thread_info::ptx_fetch_inst( inst_t &inst ) const
+{
+   addr_t pc = get_pc();
+   const ptx_instruction *pI = m_func_info->get_instruction(pc);
+   inst = (const inst_t&)*pI;
+   assert( inst.valid() );
 }
 
-void function_info::ptx_decode_inst( ptx_thread_info *thread, 
-                                     unsigned *op_type, 
-                                     int *i1, int *i2, int *i3, int *i4, 
-                                     int *o1, int *o2, int *o3, int *o4, 
-                                     int *vectorin, 
-                                     int *vectorout,
-                                     int *arch_reg,
-                                     int *pred,
-                                     int *ar1, int *ar2 )
+void ptx_instruction::pre_decode()
 {
-   addr_t pc = thread->get_pc();
-   unsigned index = pc - m_start_PC;
-   assert( index < m_instr_mem_size );
-   ptx_instruction *pI = m_instr_mem[index]; //get instruction from m_instr_mem[PC]
+   pc = m_PC;
+   isize = m_inst_size;
+   for( unsigned i=0; i<4; i++) {
+       out[i] = 0;
+       in[i] = 0;
+   }
+   is_vectorin = 0;
+   is_vectorout = 0;
+   std::fill_n(arch_reg, MAX_REG_OPERANDS, -1);
+   pred = 0;
+   ar1 = 0;
+   ar2 = 0;
 
    bool has_dst = false ;
-   int opcode = pI->get_opcode(); //determine the opcode
+   int opcode = get_opcode(); //determine the opcode
 
-   switch ( pI->get_opcode() ) {
+   switch ( get_opcode() ) {
 #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break;
 #include "opcodes.def"
 #undef OP_DEF
    default:
-      printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() );
+      printf( "Execution error: Invalid opcode (0x%x)\n", get_opcode() );
       break;
    }
 
-   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_opcode_info(this,opcode,&cycles,&op);
 
    // Get register operands
    int n=0,m=0;
-   ptx_instruction::const_iterator op=pI->op_iter_begin();
-   for ( ; op != pI->op_iter_end(); op++, n++ ) { //process operands
+   ptx_instruction::const_iterator opr=op_iter_begin();
+   for ( ; opr != op_iter_end(); opr++, n++ ) { //process operands
 
-      const operand_info &o = *op;
+      const operand_info &o = *opr;
       if ( has_dst && n==0 ) {
          if ( o.is_reg() ) { //but is destination an actual register? (seems like it fails if it's a vector)
-            *o1 = o.reg_num();
+            out[0] = o.reg_num();
             arch_reg[0] = o.arch_reg_num();
          } else if ( o.is_vector() ) { //but is destination an actual register? (seems like it fails if it's a vector)
-            *vectorin = 1;
-            *o1 = o.reg1_num();
-            *o2 = o.reg2_num();
-            *o3 = o.reg3_num();
-            *o4 = o.reg4_num();
-            for (int i = 0; i < 4; i++) 
+            is_vectorin = 1;
+            unsigned num_elem = o.get_vect_nelem();
+            if( num_elem >= 1 ) out[0] = o.reg1_num();
+            if( num_elem >= 2 ) out[1] = o.reg2_num();
+            if( num_elem >= 3 ) out[2] = o.reg3_num();
+            if( num_elem >= 4 ) out[3] = o.reg4_num();
+            for (int i = 0; i < num_elem; i++) 
                arch_reg[i] = o.arch_reg_num(i);
          }
       } else {
@@ -648,21 +659,21 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread,
             int reg_num = o.reg_num();
             arch_reg[m + 4] = o.arch_reg_num();
             switch ( m ) {
-            case 0: *i1 = reg_num; break;
-            case 1: *i2 = reg_num; break;
-            case 2: *i3 = reg_num; break;
-            default: 
-               break; 
+            case 0: in[0] = reg_num; break;
+            case 1: in[1] = reg_num; break;
+            case 2: in[2] = reg_num; break;
+            default: break; 
             }
             m++;
          } else if ( o.is_vector() ) {
             assert(m == 0); //only support 1 vector operand (for textures) right now
-            *vectorout = 1;
-            *i1 = o.reg1_num();
-            *i2 = o.reg2_num();
-            *i3 = o.reg3_num();
-            *i4 = o.reg4_num();
-            for (int i = 0; i < 4; i++) 
+            is_vectorout = 1;
+            unsigned num_elem = o.get_vect_nelem();
+            if( num_elem >= 1 ) in[0] = o.reg1_num();
+            if( num_elem >= 2 ) in[1] = o.reg2_num();
+            if( num_elem >= 3 ) in[2] = o.reg3_num();
+            if( num_elem >= 4 ) in[3] = o.reg4_num();
+            for (int i = 0; i < num_elem; i++) 
                arch_reg[i + 4] = o.arch_reg_num(i);
             m+=4;
          }
@@ -670,64 +681,34 @@ void function_info::ptx_decode_inst( ptx_thread_info *thread,
    }
 
    // 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()){
-				 *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();
+   if(has_pred()) {
+	   const operand_info &p = 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
+   if( has_memory_read() || has_memory_write() ) {
+	  ptx_instruction::const_iterator op=op_iter_begin();
+	  for ( ; op != 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();
+				 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();
+	    		 ar1 = o.reg1_num();
+	    		 ar2 = o.reg2_num();
 	    	 }
 	     }
 	  }
    }
+   m_decoded=true;
 }
 
 void function_info::add_param_name_type_size( unsigned index, std::string name, int type, size_t size )
@@ -855,7 +836,7 @@ void function_info::param_to_shared( memory_space *shared_mem, symbol_table *sym
       type_info_key::type_decode(xtype,size,tmp);
 
       // Write to shared memory - offset + 0x10
-      shared_mem->write(offset+0x10,size/8,&value,NULL,NULL);
+      shared_mem->write(offset+0x10,size/8,value.pdata,NULL,NULL);
    }
 }
 
@@ -923,49 +904,45 @@ unsigned datatype2size( unsigned data_type )
 
 unsigned g_warp_active_mask;
 
-void function_info::ptx_exec_inst( ptx_thread_info *thread, 
-                                   addr_t *addr, 
-                                   memory_space_t *space, 
-                                   unsigned *data_size,
-                                   unsigned *cycles,
-                                   dram_callback_t* callback, 
-                                   unsigned warp_active_mask  )
+void ptx_thread_info::ptx_exec_inst( inst_t &inst )
 {
+   inst.memory_op = no_memory_op;
    bool skip = false;
    int op_classification = 0;
-   addr_t pc = thread->next_instr();
-   unsigned index = pc - m_start_PC;
-   assert( index < m_instr_mem_size );
-   ptx_instruction *pI = m_instr_mem[index];
+   addr_t pc = next_instr();
+   assert( pc == inst.pc ); // make sure timing model and functional model are in sync
+   const ptx_instruction *pI = m_func_info->get_instruction(pc);
+   set_npc( pc + pI->inst_size() );
+
    try {
 
-   thread->clearRPC();
-   thread->m_last_set_operand_value.u64 = 0;
+   clearRPC();
+   m_last_set_operand_value.u64 = 0;
 
-   if(thread->is_done())
+   if(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();
-         thread->enable_debug_trace();
+   if ( g_debug_execution >= 6 || g_ptx_inst_debug_to_file) {
+      if ( (g_debug_thread_uid==0) || (get_uid() == (unsigned)g_debug_thread_uid) ) {
+         clear_modifiedregs();
+         enable_debug_trace();
       }
    }
    if( pI->has_pred() ) {
       const operand_info &pred = pI->get_pred();
-      ptx_reg_t pred_value = thread->get_operand_value(pred, pred, PRED_TYPE, thread, 0);
+      ptx_reg_t pred_value = get_operand_value(pred, pred, PRED_TYPE, this, 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;
+   g_warp_active_mask = inst.warp_active_mask;
    if( !skip ) {
       switch ( pI->get_opcode() ) {
-#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,thread); op_classification = CLASSIFICATION; break;
+#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
 #include "opcodes.def"
 #undef OP_DEF
 
@@ -976,110 +953,108 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread,
 
       // Run exit instruction if exit option included
       if(pI->is_exit())
-         exit_impl(pI,thread);
+         exit_impl(pI,this);
    }
 
    // 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();
+        (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == get_uid()) ) {
+      dim3 ctaid = get_ctaid();
+      dim3 tid = get_tid();
       fprintf(ptx_inst_debug_file,
              "[thd=%u] : (%s:%u - %s)\n",
-             thread->get_uid(),
+             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) ) {
-      dim3 ctaid = thread->get_ctaid();
-      dim3 tid = thread->get_tid();
-      printf("%u [cyc=%u][thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s)  [0x%llx]\n", 
+   if ( ptx_debug_exec_dump_cond<5>(get_uid(), pc) ) {
+      dim3 ctaid = get_ctaid();
+      dim3 tid = get_tid();
+      printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s)  [0x%llx]\n", 
              g_ptx_sim_num_insn, 
-             (unsigned)gpu_sim_cycle,
-             thread->get_uid(),
+             get_uid(),
              pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z,
-             thread->get_icount(),
+             get_icount(),
              pc, pI->source_file(), pI->source_line(), pI->get_source(),
-             thread->m_last_set_operand_value.u64 );
+             m_last_set_operand_value.u64 );
       fflush(stdout);
    }
 
    addr_t insn_memaddr = 0xFEEBDAED;
    memory_space_t insn_space = undefined_space;
+   _memory_op_t insn_memory_op = no_memory_op;
    unsigned insn_data_size = 0;
    if ( (pI->has_memory_read()  || pI->has_memory_write()) ) {
-      insn_memaddr = thread->last_eaddr();
-      insn_space = thread->last_space();
+      insn_memaddr = last_eaddr();
+      insn_space = last_space();
       unsigned to_type = pI->get_type();
       insn_data_size = datatype2size(to_type);
+      insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
    }
 
    if ( pI->get_opcode() == ATOM_OP ) {
-      insn_memaddr = thread->last_eaddr();
-      insn_space = thread->last_space();
-      callback->function = thread->last_callback().function;
-      callback->instruction = thread->last_callback().instruction;
-      callback->thread = thread;
+      insn_memaddr = last_eaddr();
+      insn_space = last_space();
+      inst.callback.function = last_callback().function;
+      inst.callback.instruction = last_callback().instruction;
+      inst.callback.thread = this;
 
       unsigned to_type = pI->get_type();
       insn_data_size = datatype2size(to_type);
    } else {
       // make sure that the callback isn't set
-      callback->function = NULL;
-      callback->instruction = NULL;
+      inst.callback.function = NULL;
+      inst.callback.instruction = NULL;
+      inst.callback.thread = NULL;
    }
 
-   // Set number of cycles for this instruction
-   int opcode = pI->get_opcode(); //determine the opcode
-   unsigned op_type;
-   get_opcode_info(pI,opcode,cycles,&op_type);
-
    if (pI->get_opcode() == TEX_OP) {
-      *addr = thread->last_eaddr();
-      *space = thread->last_space();
+      inst.memreqaddr = last_eaddr();
+      inst.space = last_space();
 
       unsigned to_type = pI->get_type();
       switch ( to_type ) {
       case B8_TYPE:
       case S8_TYPE:
       case U8_TYPE: 
-         *data_size = 1; break;
+         inst.data_size = 1; break;
       case B16_TYPE:
       case S16_TYPE:
       case U16_TYPE:
       case F16_TYPE: 
-         *data_size = 2; break;
+         inst.data_size = 2; break;
       case B32_TYPE:
       case S32_TYPE:
       case U32_TYPE:
       case F32_TYPE: 
-         *data_size = 4; break;
+         inst.data_size = 4; break;
       case B64_TYPE:
       case S64_TYPE:
       case U64_TYPE:
       case F64_TYPE: 
-         *data_size = 8; break;
+         inst.data_size = 8; break;
       default: assert(0); break;
       }
    }
 
    // Output register 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()) ) {
-      thread->dump_modifiedregs(ptx_inst_debug_file);
+        (g_ptx_inst_debug_thread_uid == 0 || g_ptx_inst_debug_thread_uid == get_uid()) ) {
+      dump_modifiedregs(ptx_inst_debug_file);
+      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(stdout);
+      if ( ptx_debug_exec_dump_cond<6>(get_uid(), pc) )
+         dump_modifiedregs(stdout);
    }
    if ( g_debug_execution >= 10 ) {
-      if ( ptx_debug_exec_dump_cond<10>(thread->get_uid(), pc) )
-         thread->dump_regs(stdout);
+      if ( ptx_debug_exec_dump_cond<10>(get_uid(), pc) )
+         dump_regs(stdout);
    }
-   thread->update_pc();
+   update_pc( pI->inst_size() );
    g_ptx_sim_num_insn++;
    ptx_file_line_stats_add_exec_count(pI);
    if ( gpgpu_ptx_instruction_classification ) {
@@ -1103,17 +1078,24 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread,
       StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int)  pI->get_opcode() );
    }
    if ( (g_ptx_sim_num_insn % 100000) == 0 ) {
-      dim3 ctaid = thread->get_ctaid();
-      dim3 tid = thread->get_tid();
+      dim3 ctaid = get_ctaid();
+      dim3 tid = get_tid();
       printf("GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n",
              g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z );
       fflush(stdout);
    }
 
    // "Return values"
-   *space = insn_space;
-   *addr = insn_memaddr;
-   *data_size = insn_data_size;
+   if(!skip) {
+      inst.space = insn_space;
+      inst.memreqaddr = insn_memaddr;
+      inst.data_size = insn_data_size;
+      inst.memory_op = insn_memory_op;
+   } else {
+      inst.memreqaddr = 0xFEEBDAED;
+      inst.space = undefined_space;
+      inst.memory_op = no_memory_op;
+   }
 
    } catch ( int x  ) {
       printf("GPGPU-Sim PTX: ERROR (%d) executing intruction (%s:%u)\n", x, pI->source_file(), pI->source_line() );
@@ -1122,10 +1104,6 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread,
    }
 }
 
-unsigned g_gx, g_gy, g_gz;
-
-dim3 g_cudaGridDim, g_cudaBlockDim;
-
 unsigned g_cta_launch_sid;
 std::list<ptx_thread_info *> g_active_threads;
 std::map<unsigned,unsigned> g_sm_idx_offset_next;
@@ -1134,45 +1112,39 @@ std::map<unsigned,memory_space*> g_shared_memory_lookup;
 std::map<unsigned,ptx_cta_info*> g_ptx_cta_lookup;
 std::map<unsigned,std::map<unsigned,memory_space*> > g_local_memory_lookup;
 
-// return number of blocks in grid
-unsigned ptx_sim_grid_size()
-{
-   return g_cudaGridDim.x * g_cudaGridDim.y * g_cudaGridDim.z;
-}
-
-void set_option_gpgpu_spread_blocks_across_cores(int option)
-{
-   gpgpu_option_spread_blocks_across_cores = option;
-}
-
 void set_param_gpgpu_num_shaders(int num_shaders)
 {
    gpgpu_param_num_shaders = num_shaders;
 }
 
-unsigned ptx_sim_cta_size()
+const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(function_info *kernel) 
 {
-   return g_cudaBlockDim.x * g_cudaBlockDim.y * g_cudaBlockDim.z;
-} 
-
-const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info() {
-   return g_entrypoint_func_info->get_kernel_info();
+   return kernel->get_kernel_info();
 }
 
-void ptx_sim_free_sm( ptx_thread_info** thread_info )
+const inst_t *ptx_fetch_inst( address_type pc )
 {
+    return function_info::pc_to_instruction(pc);
 }
 
-unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid,unsigned threads_left,unsigned num_threads, core_t *core, unsigned hw_cta_id, unsigned hw_warp_id )
+unsigned ptx_sim_init_thread( kernel_info_t &kernel,
+                              ptx_thread_info** thread_info,
+                              int sid,
+                              unsigned tid,
+                              unsigned threads_left,
+                              unsigned num_threads, 
+                              core_t *core, 
+                              unsigned hw_cta_id, 
+                              unsigned hw_warp_id )
 {
    if ( *thread_info != NULL ) {
       ptx_thread_info *thd = *thread_info;
       assert( thd->is_done() );
       if ( g_debug_execution==-1 ) {
          dim3 ctaid = thd->get_ctaid();
-         dim3 tid = thd->get_tid();
+         dim3 t = thd->get_tid();
          printf("GPGPU-Sim PTX simulator:  thread exiting ctaid=(%u,%u,%u) tid=(%u,%u,%u) uid=%u\n",
-                ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, thd->get_uid() );
+                ctaid.x,ctaid.y,ctaid.z,t.x,t.y,t.z, thd->get_uid() );
          fflush(stdout);
       }
       thd->m_cta_info->assert_barrier_empty();
@@ -1197,11 +1169,11 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid
       return 1;
    }
 
-   if ( g_gx >= g_cudaGridDim.x  || g_gy >= g_cudaGridDim.y || g_gz >= g_cudaGridDim.z ) {
+   if ( kernel.no_more_ctas_to_run() ) {
       return 0; //finished!
    }
 
-   if ( threads_left < ptx_sim_cta_size() ) {
+   if ( threads_left < kernel.threads_per_cta() ) {
       return 0;
    }
 
@@ -1214,25 +1186,12 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid
    ptx_cta_info *cta_info = NULL;
    memory_space *shared_mem = NULL;
 
-   unsigned cta_size = ptx_sim_cta_size(); //blocksize
+   unsigned cta_size = kernel.threads_per_cta();
    unsigned sm_offset = g_sm_idx_offset_next[sid];
    unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 
    assert( max_cta_per_sm > 0 );
 
-   unsigned sm_idx = sid*max_cta_per_sm + sm_offset;
-   sm_idx = max_cta_per_sm*sid + tid/cta_size;
-
-   if (!gpgpu_option_spread_blocks_across_cores) {
-      // update offset...
-      if ( (sm_offset + 1) >= max_cta_per_sm ) {
-         sm_offset = 0;
-      } else {
-         sm_offset++;
-      }
-      g_sm_idx_offset_next[sid] = sm_offset;
-   } else {
-      sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid;
-   }
+   unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid;
 
    if ( g_shared_memory_lookup.find(sm_idx) == g_shared_memory_lookup.end() ) {
       if ( g_debug_execution >= 1 ) {
@@ -1256,67 +1215,57 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid
    }
 
    std::map<unsigned,memory_space*> &local_mem_lookup = g_local_memory_lookup[sid];
-   unsigned new_tid;
-   for ( unsigned tz=0; tz < g_cudaBlockDim.z; tz++ ) {
-      for ( unsigned ty=0; ty < g_cudaBlockDim.y; ty++ ) {
-         for ( unsigned tx=0; tx < g_cudaBlockDim.x; tx++ ) {
-            new_tid = tx + g_cudaBlockDim.x*ty + g_cudaBlockDim.x*g_cudaBlockDim.y*tz;
-            new_tid += tid;
-            ptx_thread_info *thd = new ptx_thread_info();
-
-            memory_space *local_mem = NULL;
-            std::map<unsigned,memory_space*>::iterator l = local_mem_lookup.find(new_tid);
-            if ( l != local_mem_lookup.end() ) {
-               local_mem = l->second;
-            } else {
-               char buf[512];
-               snprintf(buf,512,"local_%u_%u", sid, new_tid);
-               local_mem = new memory_space_impl<32>(buf,32);
-               local_mem_lookup[new_tid] = local_mem;
-            }
-            thd->set_info(g_entrypoint_func_info);
-            thd->set_nctaid(g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z);
-            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);
-            thd->set_core(NULL);
-            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;
-            if ( g_debug_execution==-1 ) {
-               printf("GPGPU-Sim PTX simulator:  allocating thread ctaid=(%u,%u,%u) tid=(%u,%u,%u) @ 0x%Lx\n",
-                      g_gx,g_gy,g_gz,tx,ty,tz, (unsigned long long)thd );
-               fflush(stdout);
-            }
-            g_active_threads.push_back(thd);
-         }
+   while( kernel.more_threads_in_cta() ) {
+      dim3 ctaid3d = kernel.get_next_cta_id();
+      unsigned new_tid = kernel.get_next_thread_id();
+      dim3 tid3d = kernel.get_next_thread_id_3d();
+      kernel.increment_thread_id();
+      new_tid += tid;
+      ptx_thread_info *thd = new ptx_thread_info();
+   
+      memory_space *local_mem = NULL;
+      std::map<unsigned,memory_space*>::iterator l = local_mem_lookup.find(new_tid);
+      if ( l != local_mem_lookup.end() ) {
+         local_mem = l->second;
+      } else {
+         char buf[512];
+         snprintf(buf,512,"local_%u_%u", sid, new_tid);
+         local_mem = new memory_space_impl<32>(buf,32);
+         local_mem_lookup[new_tid] = local_mem;
       }
+      thd->set_info(kernel.entry());
+      thd->set_nctaid(kernel.get_grid_dim());
+      thd->set_ntid(kernel.get_cta_dim());
+      thd->set_ctaid(ctaid3d);
+      thd->set_tid(tid3d);
+      if( kernel.entry()->get_ptx_version().extensions() ) 
+         thd->cpy_tid_to_reg(tid3d);
+      thd->set_hw_tid((unsigned)-1);
+      thd->set_hw_wid((unsigned)-1);
+      thd->set_hw_ctaid((unsigned)-1);
+      thd->set_core(NULL);
+      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;
+      if ( g_debug_execution==-1 ) {
+         printf("GPGPU-Sim PTX simulator:  allocating thread ctaid=(%u,%u,%u) tid=(%u,%u,%u) @ 0x%Lx\n",
+                ctaid3d.x,ctaid3d.y,ctaid3d.z,tid3d.x,tid3d.y,tid3d.z, (unsigned long long)thd );
+         fflush(stdout);
+      }
+      g_active_threads.push_back(thd);
    }
    if ( g_debug_execution==-1 ) {
       printf("GPGPU-Sim PTX simulator:  <-- FINISHING THREAD ALLOCATION\n");
       fflush(stdout);
    }
 
-   g_gx++;
-   if ( g_gx >= g_cudaGridDim.x ) {
-      g_gx = 0;
-      g_gy++;
-      if ( g_gy >= g_cudaGridDim.y ) {
-         g_gy = 0;
-         g_gz++;
-      }
-   }
+   kernel.increment_cta_id();
 
    g_cta_launch_sid = -1;
 
@@ -1336,7 +1285,7 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid
 
 void init_inst_classification_stat() {
    char kernelname[256] ="";
-#define MAX_CLASS_KER 256
+#define MAX_CLASS_KER 1024
    if (!g_inst_classification_stat) g_inst_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*));
    snprintf(kernelname, MAX_CLASS_KER, "Kernel %d Classification\n",g_ptx_kernel_count  );         
    assert( g_ptx_kernel_count < MAX_CLASS_KER ) ; // a static limit on number of kernels increase it if it fails! 
@@ -1350,21 +1299,12 @@ void init_inst_classification_stat() {
 std::map<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL;
 std::map<const void*,std::string> *g_host_to_kernel_entrypoint_name_lookup=NULL;
 
-void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args,
-                                         struct dim3 gridDim, struct dim3 blockDim ) 
+function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled )
 {
-   g_gx=0;
-   g_gy=0;
-   g_gz=0;
-   g_cudaGridDim = gridDim;
-   g_cudaBlockDim = blockDim;
-   g_sm_idx_offset_next.clear();
-   g_sm_next_index = 0;  
-
    if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) ==
         g_host_to_kernel_entrypoint_name_lookup->end() ) {
-      printf("GPGPU-Sim PTX: ERROR ** cannot locate PTX entry point\n" );
-      printf("GPGPU-Sim PTX: existing entry points: \n");
+      printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" );
+      printf("GPGPU-Sim PTX: registered PTX kernels: \n");
       std::map<const void*,std::string>::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin();
       for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) {
          printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str());
@@ -1372,71 +1312,77 @@ void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_
       printf("\n");
       abort();
    } 
-
-   std::string kname = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key];
-   printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n",
-          kname.c_str(), g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z,g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z, 
-          g_ptx_thread_info_uid_next );
-
-   if ( g_kernel_name_to_function_lookup->find(kname) ==
+   kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key];
+   if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) ==
         g_kernel_name_to_function_lookup->end() ) {
-      printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() );
+      printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", 
+             kernel_func_name_mangled.c_str() );
       abort();
    }
-   g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname];
+   return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled];
+}
 
-   unsigned argcount=0;
-   struct gpgpu_ptx_sim_arg *tmparg = args;
-   while (tmparg) {
-      tmparg = tmparg->m_next;
-      argcount++;
-   }
+const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key)
+{
+   std::string kname;
+   function_info *finfo = get_kernel(kernel_key,kname);
+   return finfo->get_kernel_info();
+}
+
+size_t get_kernel_code_size( class function_info *entry )
+{
+   return entry->get_function_size();
+}
 
+kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args,
+                                         struct dim3 gridDim, struct dim3 blockDim ) 
+{
+   g_sm_idx_offset_next.clear();
+   g_sm_next_index = 0;  
+   std::string kname;
+   function_info *entry = get_kernel(kernel_key,kname);
+
+   printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n",
+          kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z, 
+          g_ptx_thread_info_uid_next );
+
+
+   unsigned argcount=args.size();
    unsigned argn=1;
-   while (args) {
-      g_func_info->add_param_data(argcount-argn,args);
-      args = args->m_next;
+   for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
+      entry->add_param_data(argcount-argn,&(*a));
       argn++;
    }
-   g_func_info->finalize(g_param_mem);
+
+   entry->finalize(g_param_mem);
    g_ptx_kernel_count++; 
    if ( gpgpu_ptx_instruction_classification ) {
       init_inst_classification_stat();
    }
    fflush(stdout);
+
+   return kernel_info_t(gridDim,blockDim,entry);
 }
 
-void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim )
+kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim )
 {
-   g_gx=0;
-   g_gy=0;
-   g_gz=0;
-   g_cudaGridDim = gridDim;
-   g_cudaBlockDim = blockDim;
    g_sm_idx_offset_next.clear();
    g_sm_next_index = 0;  
 
-   g_entrypoint_func_info = g_func_info = entry;
-
-   unsigned argcount=0;
-   struct gpgpu_ptx_sim_arg *tmparg = args;
-   while (tmparg) {
-      tmparg = tmparg->m_next;
-      argcount++;
-   }
-
+   unsigned argcount=args.size();
    unsigned argn=1;
-   while (args) {
-      g_func_info->add_param_data(argcount-argn,args);
-      args = args->m_next;
+   for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
+      entry->add_param_data(argcount-argn,&(*a));
       argn++;
    }
-   g_func_info->finalize(g_param_mem);
+   entry->finalize(g_param_mem);
    g_ptx_kernel_count++; 
    if ( gpgpu_ptx_instruction_classification ) {
       init_inst_classification_stat();
    }
    fflush(stdout);
+
+   return kernel_info_t(gridDim,blockDim,entry);
 }
 
 const char *g_gpgpusim_version_string = "2.1.1b (beta)";
@@ -1625,7 +1571,7 @@ ptx_cta_info *g_func_cta_info = NULL;
 
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
-void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args)
+void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args)
 {
    printf("GPGPU-Sim: Performing Functional Simulation...\n");
 
@@ -1634,12 +1580,11 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl
    exit(1);
    
    time_t end_time, elapsed_time, days, hrs, minutes, sec;
-   int i1, i2, i3, i4, o1, o2, o3, o4;
-   int vectorin, vectorout;
-   int pred;
-   int ar1, ar2;
 
-   gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim);
+   kernel_info_t kernel = gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim);
+
+   std::string kname;
+   function_info *finfo = get_kernel(kernel_key,kname);
 
    memory_space *shared_mem = new memory_space_impl<16*1024>("shared",4);
 
@@ -1647,46 +1592,40 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl
 
    if ( g_func_cta_info == NULL )
       g_func_cta_info = new ptx_cta_info(0);
-
-   for ( unsigned gx=0; gx < gridDim.x; gx++ ) {
-      for ( unsigned gy=0; gy < gridDim.y; gy++ ) {
-         for ( unsigned gz=0; gz < gridDim.z; gz++ ) {
+   while( !kernel.no_more_ctas_to_run() ) {
             std::list<ptx_thread_info *> active_threads;
             std::list<ptx_thread_info *> blocked_threads;
+            dim3 ctaid3d = kernel.get_next_cta_id();
+            kernel.increment_cta_id();
 
             g_func_cta_info->check_cta_thread_status_and_reset();
-
-            for ( unsigned tx=0; tx < blockDim.x; tx++ ) {
-               for ( unsigned ty=0; ty < blockDim.y; ty++ ) {
-                  for ( unsigned tz=0; tz < blockDim.z; tz++ ) {
+            while( kernel.more_threads_in_cta() ) {
                      memory_space *local_mem = NULL;
                      ptx_thread_info *thd = new ptx_thread_info();
-
-                     unsigned lm_idx = blockDim.x*blockDim.y*tz + blockDim.x * ty + tx;
+                     dim3 tid3d = kernel.get_next_thread_id_3d();
+                     unsigned lm_idx = kernel.get_next_thread_id();
+                     kernel.increment_thread_id();
                      std::map<unsigned,memory_space*>::iterator lm=lm_lookup.find(lm_idx);
                      if ( lm == lm_lookup.end() ) {
                         char buf[1024];
-                        snprintf(buf,1024,"local_(%u,%u,%u)", tx, ty, tz );
+                        snprintf(buf,1024,"local_(%u,%u,%u)", tid3d.x, tid3d.y, tid3d.z );
                         local_mem = new memory_space_impl<32>(buf,32);
                         lm_lookup[lm_idx] = local_mem;
                      } else {
                         local_mem = lm->second;
                      }
 
-
-                     thd->set_info(g_func_info);
-                     thd->set_nctaid(gridDim.x,gridDim.y,gridDim.z);
-                     thd->set_ntid(blockDim.x, blockDim.y, blockDim.z);
-                     thd->set_ctaid(gx,gy,gz);
-                     thd->set_tid(tx,ty,tz);
+                     thd->set_info(finfo);
+                     thd->set_nctaid(ctaid3d);
+                     thd->set_ntid(kernel.get_cta_dim());
+                     thd->set_ctaid(ctaid3d);
+                     thd->set_tid(tid3d);
                      thd->set_valid();
                      thd->m_shared_mem = shared_mem;
                      thd->m_local_mem = local_mem;
                      thd->m_cta_info = g_func_cta_info;
                      g_func_cta_info->add_thread(thd);
                      active_threads.push_back(thd);
-                  }
-               }
             }
 
             while ( !(active_threads.empty() && blocked_threads.empty()) ) {
@@ -1720,22 +1659,12 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl
                      break;
                   }
 
-                  unsigned op_type;
-                  addr_t addr;
-                  unsigned cycles;
-                  memory_space_t space;
-                  int arch_reg[MAX_REG_OPERANDS] = { -1 };
-                  unsigned data_size;
-                  dram_callback_t callback;
-                  unsigned warp_active_mask = (unsigned)-1; // vote instruction with diverged warps won't execute correctly
-                                                            // in functional simulation mode
-
-                  g_func_info->ptx_decode_inst( thread, &op_type, &i1, &i2, &i3, &i4, &o1, &o2, &o3, &o4, &vectorin, &vectorout, arch_reg, &pred, &ar1, &ar2 );
-                  g_func_info->ptx_exec_inst( thread, &addr, &space, &data_size, &cycles, &callback, warp_active_mask );
+                  inst_t inst;
+                  inst.warp_active_mask = (unsigned)-1; // vote instruction with diverged warps won't execute correctly
+                                                        // in functional simulation mode
+                  thread->ptx_exec_inst( inst );
                }
             }
-         }
-      }
    }
    printf( "GPGPU-Sim: Done functional simulation (%u instructions simulated).\n", g_ptx_sim_num_insn );
    if ( gpgpu_ptx_instruction_classification ) {
@@ -1757,70 +1686,12 @@ void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 bl
    fflush(stdout); 
 }
 
-void ptx_decode_inst( void *thd, unsigned *op, int *i1, int *i2, int *i3, int *i4, int *o1, int *o2, int *o3, int *o4, int *vectorin, int *vectorout, int *arch_reg, int *pred, int *ar1, int *ar2 )
-{
-   *op = NO_OP;
-   *o1 = 0;
-   *o2 = 0;
-   *o3 = 0;
-   *o4 = 0;
-   *i1 = 0;
-   *i2 = 0;
-   *i3 = 0;
-   *i4 = 0;
-   *vectorin = 0;
-   *vectorout = 0;
-   std::fill_n(arch_reg, MAX_REG_OPERANDS, -1);
-   *pred = 0;
-   *ar1 = 0;
-   *ar2 = 0;
-
-   if ( thd == NULL )
-      return;
-
-   ptx_thread_info *thread = (ptx_thread_info *) thd;
-   g_func_info = thread->func_info();
-   g_func_info->ptx_decode_inst(thread,op,i1,i2,i3,i4,o1,o2,o3,o4,vectorin,vectorout,arch_reg,pred,ar1,ar2);
-}
-
-extern "C" unsigned ptx_get_inst_op( void *thd)
-{
-   if ( thd == NULL )
-      return NO_OP;
-
-   ptx_thread_info *thread = (ptx_thread_info *) thd;
-   return(thread->func_info())->ptx_get_inst_op(thread);
-}
-
-void ptx_exec_inst( void *thd, address_type *addr, memory_space_t *space, unsigned *data_size, unsigned *cycles, dram_callback_t* callback, unsigned warp_active_mask )
-{
-   if ( thd == NULL )
-      return;
-   *cycles = 1;
-   ptx_thread_info *thread = (ptx_thread_info *) thd;
-   g_func_info = thread->func_info();
-   g_func_info->ptx_exec_inst( thread, addr, space, data_size, cycles, callback, warp_active_mask );
-}
-
-void ptx_dump_regs( void *thd )
-{
-   if ( thd == NULL )
-      return;
-   ptx_thread_info *t = (ptx_thread_info *) thd;
-   t->dump_regs(stdout);
-}
-
 unsigned ptx_set_tex_cache_linesize(unsigned linesize)
 {
    g_texcache_linesize = linesize;
    return 0;
 }
 
-unsigned ptx_kernel_program_size()
-{
-   return g_func_info->get_function_size();
-}
-
 unsigned translate_pc_to_ptxlineno(unsigned pc)
 {
    // this function assumes that the kernel fits inside a single PTX file
@@ -1837,12 +1708,11 @@ int g_ptxinfo_error_detected;
 static char *g_ptxinfo_kname = NULL;
 static struct gpgpu_ptx_sim_kernel_info g_ptxinfo_kinfo;
 
+static void clear_ptxinfo();
+
 extern "C" void ptxinfo_function(const char *fname )
 {
-    g_ptxinfo_kinfo.regs=0;
-    g_ptxinfo_kinfo.lmem=0;
-    g_ptxinfo_kinfo.smem=0;
-    g_ptxinfo_kinfo.cmem=0;
+    clear_ptxinfo();
     g_ptxinfo_kname = strdup(fname);
 }
 
@@ -1874,6 +1744,8 @@ void clear_ptxinfo()
     g_ptxinfo_kinfo.lmem=0;
     g_ptxinfo_kinfo.smem=0;
     g_ptxinfo_kinfo.cmem=0;
+    g_ptxinfo_kinfo.ptx_version=0;
+    g_ptxinfo_kinfo.sm_target=0;
 }
 
 void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels )
@@ -1991,14 +1863,9 @@ unsigned int get_converge_point( unsigned int pc, void *thd )
    abort(); // returning garbage!
 }
 
-void find_reconvergence_points()
-{
-    find_reconvergence_points(g_func_info);
-}
-
-void dwf_process_reconv_pts()
+void dwf_process_reconv_pts(function_info *entry)
 {
-   rec_pts tmp = find_reconvergence_points(g_func_info);
+   rec_pts tmp = find_reconvergence_points(entry);
    for (int i = 0; i < tmp.s_num_recon; ++i) {
       dwf_insert_reconv_pt(tmp.s_kernel_recon_points[i].target_pc);
    }
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index d7d3f8a..57f2837 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -23,15 +23,15 @@ extern int g_ptx_kernel_count; // used for classification stat collection purpos
 extern FILE* ptx_inst_debug_file;
 
 
-extern void   gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
-                                            struct gpgpu_ptx_sim_arg *args, 
+extern class kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
+                                            gpgpu_ptx_sim_arg_list_t args, 
                                             struct dim3 gridDim, 
                                             struct dim3 blockDim );
-extern void   gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
-                                            struct gpgpu_ptx_sim_arg *args, 
+extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+                                            gpgpu_ptx_sim_arg_list_t args, 
                                             struct dim3 gridDim, 
                                             struct dim3 blockDim );
-extern void   gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *);
+extern void   gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t );
 extern void   print_splash();
 extern void*  gpgpu_ptx_sim_malloc( size_t count );
 extern void*  gpgpu_ptx_sim_mallocarray( size_t count );
@@ -55,8 +55,8 @@ extern void read_sim_environment_variables();
 extern void register_function_implementation( const char *name, function_info *impl );
 extern void ptxinfo_cuda_addinfo();
 extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
-extern void ptx_sim_free_sm( class ptx_thread_info** thread_info );
-unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info,
+unsigned ptx_sim_init_thread( kernel_info_t &kernel,
+                              class ptx_thread_info** thread_info,
                               int sid,
                               unsigned tid,
                               unsigned threads_left,
@@ -64,12 +64,9 @@ unsigned ptx_sim_init_thread( class ptx_thread_info** thread_info,
                               class core_t *core, 
                               unsigned hw_cta_id, 
                               unsigned hw_warp_id );
-unsigned ptx_sim_cta_size();
-const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info();
-void set_option_gpgpu_spread_blocks_across_cores(int option);
-int ptx_thread_done( void *thd );
+const inst_t *ptx_fetch_inst( address_type pc );
+const struct gpgpu_ptx_sim_kernel_info* ptx_sim_kernel_info(class function_info *kernel);
 unsigned ptx_thread_donecycle( void *thr );
-int ptx_thread_get_next_pc( void *thd );
 void* ptx_thread_get_next_finfo( void *thd );
 int ptx_thread_at_barrier( void *thd );
 int ptx_thread_all_at_barrier( void *thd );
@@ -78,28 +75,10 @@ void ptx_thread_reset_barrier( void *thd );
 void ptx_thread_release_barrier( void *thd );
 void ptx_print_insn( address_type pc, FILE *fp );
 unsigned int ptx_set_tex_cache_linesize( unsigned linesize);
-void ptx_decode_inst( void *thd, 
-                      unsigned *op, 
-                      int *i1, 
-                      int *i2, 
-                      int *i3, 
-                      int *i4, 
-                      int *o1, 
-                      int *o2, 
-                      int *o3, 
-                      int *o4, 
-                      int *vectorin, 
-                      int *vectorout, 
-                      int *arch_reg, 
-                      int *pred,
-                      int *ar1, 
-                      int *ar2 );
-void ptx_exec_inst( void *thd, 
-                    address_type *addr, 
-                    memory_space_t *space, 
-                    unsigned *data_size,
-                    unsigned *cycles, 
-                    dram_callback_t* callback, 
-                    unsigned warp_active_mask );
+
+function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled );
+void dwf_process_reconv_pts(function_info *entry);
+void set_param_gpgpu_num_shaders(int num_shaders);
+unsigned int get_converge_point(unsigned int pc, void *thd);
 
 #endif
diff --git a/src/cuda-sim/dram_callback.h b/src/cuda-sim/dram_callback.h
index 9336ee0..fee97eb 100644
--- a/src/cuda-sim/dram_callback.h
+++ b/src/cuda-sim/dram_callback.h
@@ -75,10 +75,4 @@
 #ifndef __DRAM_CALLBACK_H__
 #define __DRAM_CALLBACK_H__
 
-typedef struct {
-   void (*function)(void* pI, void* gOldGThread);
-   void* instruction;
-   void* thread;// callback has to abuse g_thread when executed
-}dram_callback_t;
-
 #endif // #ifndef __DRAM_CALLBACK_H__
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index d9a9630..78996bb 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -651,16 +651,20 @@ void ptx_thread_info::set_vector_operand_values( const operand_info &dst,
                                                  const ptx_reg_t &data1, 
                                                  const ptx_reg_t &data2, 
                                                  const ptx_reg_t &data3, 
-                                                 const ptx_reg_t &data4, 
-                                                 unsigned num_elements )
+                                                 const ptx_reg_t &data4 )
 {
-   set_reg(dst.vec_symbol(0), data1);
-   set_reg(dst.vec_symbol(1), data2);
-   if (num_elements > 2) {
-      set_reg(dst.vec_symbol(2), data3);
-      if (num_elements > 3) {
-         set_reg(dst.vec_symbol(3), data4);
-      }
+   unsigned num_elements = dst.get_vect_nelem(); 
+   if (num_elements > 0) {
+       set_reg(dst.vec_symbol(0), data1);
+       if (num_elements > 1) {
+           set_reg(dst.vec_symbol(1), data2);
+           if (num_elements > 2) {
+              set_reg(dst.vec_symbol(2), data3);
+              if (num_elements > 3) {
+                 set_reg(dst.vec_symbol(3), data4);
+              }
+           }
+       }
    }
 
    m_last_set_operand_value = data1;
@@ -1084,6 +1088,8 @@ void atom_callback( void* ptx_inst, void* thd )
    // Write operation result into global memory
    // (i.e. copy src1_data to dst)
    g_global_mem->write(src1_data.u32,size/8,&op_result.s64,thread,pI);
+   gpgpu_sim *gpu = thread->get_gpu();
+   gpu->decrement_atomic_count(thread->get_hw_sid(),thread->get_hw_wid());
 }
 
 // atom_impl will now result in a callback being called in mem_ctrl_pop (gpu-sim.c)
@@ -1113,7 +1119,7 @@ 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, dst, U32_TYPE, thread, 1);
-   assert( b.u32 == 0 ); // not clear what should happen if this is not zero
+   assert( b.u32 == 0 ); // support for bar.sync a{,b}; where a != 0 not yet implemented
 }
 
 void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
@@ -1194,20 +1200,80 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread )
 
    unsigned sid = thread->get_hw_sid();
    unsigned tid = thread->get_hw_tid();
+   gpgpu_sim *gpu = thread->get_gpu();
    unsigned callee_pc=0, callee_rpc=0;
-   if( gpgpu_simd_model == POST_DOMINATOR ) {
-      get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc);
+   if( gpu->simd_model() == POST_DOMINATOR ) {
+      gpu->get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc);
       assert( callee_pc == thread->get_pc() );
    }
 
-   thread->callstack_push(callee_pc+1,callee_rpc,return_var_src,return_var_dst,call_uid_next++);
+   thread->callstack_push(callee_pc + pI->inst_size(), callee_rpc, return_var_src, return_var_dst, call_uid_next++);
 
    copy_buffer_list_into_frame(thread, arg_values);
 
    thread->set_npc(target_func);
 }
 
-void clz_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
+//Ptxplus version of call instruction. Jumps to a label not a different Kernel.
+void callp_impl( const ptx_instruction *pI, ptx_thread_info *thread )
+{
+   
+   static unsigned call_uid_next = 1;
+
+   const operand_info &target  = pI->dst();
+   ptx_reg_t target_pc = thread->get_operand_value(target, target, U32_TYPE, thread, 1);
+
+   const symbol *return_var_src = NULL;
+   const symbol *return_var_dst = NULL;
+
+   unsigned sid = thread->get_hw_sid();
+   unsigned tid = thread->get_hw_tid();
+   gpgpu_sim *gpu = thread->get_gpu();
+   unsigned callee_pc=0, callee_rpc=0;
+   if( gpu->simd_model() == POST_DOMINATOR ) {
+      gpu->get_pdom_stack_top_info(sid,tid,&callee_pc,&callee_rpc);
+      assert( callee_pc == thread->get_pc() );
+   } 
+
+   thread->callstack_push_plus(callee_pc + pI->inst_size(), callee_rpc, return_var_src, return_var_dst, call_uid_next++);
+   thread->set_npc(target_pc);
+}
+
+void clz_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();
+
+   unsigned i_type = pI->get_type();
+   a = thread->get_operand_value(src1, dst, i_type, thread, 1);
+
+   int max;
+   unsigned long long mask;
+   d.u64 = 0;
+
+   switch ( i_type ) {
+   case B32_TYPE:
+      max = 32;
+      mask = 0x80000000;
+      break;
+   case B64_TYPE:
+      max = 64;
+      mask = 0x8000000000000000;
+      break;
+   default:
+      printf("Execution error: type mismatch with instruction\n");
+      assert(0);
+      break;
+   }
+
+   while ((d.u32 < max) && ((a.u64&mask) == 0) ) {
+      d.u32++;
+      a.u64 = a.u64 << 1;
+   }
+
+   thread->set_operand_value(dst,d, B32_TYPE, thread, pI);
+}
 
 void cnot_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
@@ -1313,7 +1379,7 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign,
 
    ptx_reg_t y;
    if ( to_sign == 1 ) { // convert to 64-bit number first?
-      int tmp = cuda_math::__internal_float2int(x.f32, mode);
+      int tmp = cuda_math::float2int(x.f32, mode);
       if ((x.u32 & 0x7f800000) == 0)
          tmp = 0; // round denorm. FP to 0
       if (saturation_mode && to_width < 32) {
@@ -1327,7 +1393,7 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign,
       default: assert(0); break;
       }
    } else if ( to_sign == 0 ) {
-      unsigned int tmp = cuda_math::__internal_float2uint(x.f32, mode);
+      unsigned int tmp = cuda_math::float2uint(x.f32, mode);
       if ((x.u32 & 0x7f800000) == 0)
          tmp = 0; // round denorm. FP to 0
       if (saturation_mode && to_width < 32) {
@@ -1432,10 +1498,10 @@ ptx_reg_t s2f( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign,
       case 16: assert(0); break;
       case 32: 
          switch (rounding_mode) {
-         case RZ_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; 
+         case RZ_OPTION: y.f32 = cuda_math::__ll2float_rz(y.s64); break; 
          case RN_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break;
-         case RM_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break; 
-         case RP_OPTION: y.f32 = cuda_math::__ll2float_rn(y.s64); break;
+         case RM_OPTION: y.f32 = cuda_math::__ll2float_rd(y.s64); break; 
+         case RP_OPTION: y.f32 = cuda_math::__ll2float_ru(y.s64); break;
          default: break; 
          }
          break;
@@ -1549,9 +1615,9 @@ ptx_reg_t d2d( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign,
       break;          
    case RNI_OPTION: 
 #if CUDART_VERSION >= 3000
-      y.f64 = nearbyint(x.f32); 
+      y.f64 = nearbyint(x.f64); 
 #else
-      y.f64 = cuda_math::__internal_nearbyint(x.f64); 
+      y.f64 = cuda_math::__internal_nearbyintf(x.f64); 
 #endif
       break;          
    case RMI_OPTION: 
@@ -1816,7 +1882,8 @@ void cvta_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    } else {
       switch( space.get_type() ) {
       case shared_space: to_addr_hw = shared_to_generic( smid, from_addr_hw ); break;
-      case local_space:  to_addr_hw =  local_to_generic( smid, hwtid, from_addr_hw ); break;
+      case local_space:  to_addr_hw =  local_to_generic( smid, hwtid, from_addr_hw )
+                                      + thread->get_local_mem_stack_pointer(); break; // add stack ptr here so that it can be passed as a pointer at function call 
       case global_space: to_addr_hw = global_to_generic( from_addr_hw ); break;
       default: abort();
       }
@@ -2018,11 +2085,11 @@ void ld_exec( const ptx_instruction *pI, ptx_thread_info *thread )
          mem->read(addr+2*size/8,size/8,&data3.s64);
          if (vector_spec != V3_TYPE) { //v4
             mem->read(addr+3*size/8,size/8,&data4.s64);
-            thread->set_vector_operand_values(dst,data1,data2,data3,data4, 4);
+            thread->set_vector_operand_values(dst,data1,data2,data3,data4);
          } else //v3
-            thread->set_vector_operand_values(dst,data1,data2,data3,data3,3);
+            thread->set_vector_operand_values(dst,data1,data2,data3,data3);
       } else //v2
-         thread->set_vector_operand_values(dst,data1,data2,data2,data2,2);
+         thread->set_vector_operand_values(dst,data1,data2,data2,data2);
    }
    thread->m_last_effective_address = addr;
    thread->m_last_memory_space = space; 
@@ -2249,7 +2316,10 @@ void max_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    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); }
+void membar_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
+{ 
+   // handled by timing simulator 
+}
 
 void min_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
@@ -2332,16 +2402,16 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread )
          unsigned bits_per_dst_elem = nbits_to_move / nelem;
          for( unsigned i=0; i < nelem; i++ ) {
             switch(bits_per_dst_elem) {
-            case 8:  v[i].u8  = tmp_bits.u64  & (((unsigned long long) 0xFF) << (8*i)); break;
-            case 16: v[i].u16 = tmp_bits.u64  & (((unsigned long long) 0xFFFF) << (16*i)); break;
-            case 32: v[i].u32 = tmp_bits.u64  & (((unsigned long long) 0xFFFFFFFF) << (32*i)); break;
+            case 8:  v[i].u8  = (tmp_bits.u64 >> (8*i)) & ((unsigned long long) 0xFF); break;
+            case 16: v[i].u16 = (tmp_bits.u64 >> (16*i)) & ((unsigned long long) 0xFFFF); break;
+            case 32: v[i].u32 = (tmp_bits.u64 >> (32*i)) & ((unsigned long long) 0xFFFFFFFF); break;
             default:
                printf("Execution error: mov pack/unpack with unsupported source/dst size ratio (dst)\n");
                assert(0);
                break;
             }
          }
-         thread->set_vector_operand_values(dst,v[0],v[1],v[2],v[3],nelem);
+         thread->set_vector_operand_values(dst,v[0],v[1],v[2],v[3]);
       } else {
          thread->set_operand_value(dst,tmp_bits, i_type, thread, pI);
       }
@@ -2707,6 +2777,19 @@ void ret_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    }
 }
 
+//Ptxplus version of ret instruction.
+void retp_impl( const ptx_instruction *pI, ptx_thread_info *thread )
+{
+   bool empty = thread->callstack_pop_plus();
+   if( empty ) {
+      core_t *sc = thread->get_core();
+      unsigned warp_id = thread->get_hw_wid();
+      sc->warp_exit(warp_id);
+      thread->m_cta_info->register_thread_exit(thread);
+      thread->set_done();
+   }
+}
+
 void rsqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 { 
    ptx_reg_t a, d;
@@ -3689,7 +3772,7 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread )
       assert(0);
    }
    thread->m_last_memory_space = tex_space; 
-   thread->set_vector_operand_values(dst,data1,data2,data3,data4,4);
+   thread->set_vector_operand_values(dst,data1,data2,data3,data4);
 }
 
 void txq_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
@@ -3792,10 +3875,6 @@ void inst_not_implemented( const ptx_instruction * pI )
    abort();
 }
 
-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;
diff --git a/src/cuda-sim/memory.cc b/src/cuda-sim/memory.cc
index b825f12..aecfd31 100644
--- a/src/cuda-sim/memory.cc
+++ b/src/cuda-sim/memory.cc
@@ -83,10 +83,35 @@ template<unsigned BSIZE> memory_space_impl<BSIZE>::memory_space_impl( std::strin
 template<unsigned BSIZE> void memory_space_impl<BSIZE>::write( mem_addr_t addr, size_t length, const void *data, class ptx_thread_info *thd, const ptx_instruction *pI)
 {
    mem_addr_t index = addr >> m_log2_block_size;
-   unsigned offset = addr & (BSIZE-1);
-   unsigned nbytes = length;
-   assert( (addr+length) <= (index+1)*BSIZE );
-   m_data[index].write(offset,nbytes,(const unsigned char*)data);
+   if ( (addr+length) <= (index+1)*BSIZE ) {
+      // fast route for intra-block access 
+      unsigned offset = addr & (BSIZE-1);
+      unsigned nbytes = length;
+      m_data[index].write(offset,nbytes,(const unsigned char*)data);
+   } else {
+      // slow route for inter-block access
+      unsigned nbytes_remain = length;
+      unsigned src_offset = 0; 
+      mem_addr_t current_addr = addr; 
+
+      while (nbytes_remain > 0) {
+         unsigned offset = current_addr & (BSIZE-1);
+         mem_addr_t page = current_addr >> m_log2_block_size; 
+         mem_addr_t access_limit = offset + nbytes_remain; 
+         if (access_limit > BSIZE) {
+            access_limit = BSIZE;
+         } 
+         
+         size_t tx_bytes = access_limit - offset; 
+         m_data[page].write(offset, tx_bytes, &((const unsigned char*)data)[src_offset]);
+
+         // advance pointers 
+         src_offset += tx_bytes; 
+         current_addr += tx_bytes; 
+         nbytes_remain -= tx_bytes; 
+      }
+      assert(nbytes_remain == 0); 
+   }
    if( !m_watchpoints.empty() ) {
       std::map<unsigned,mem_addr_t>::iterator i;
       for( i=m_watchpoints.begin(); i!=m_watchpoints.end(); i++ ) {
@@ -97,28 +122,59 @@ template<unsigned BSIZE> void memory_space_impl<BSIZE>::write( mem_addr_t addr,
    }
 }
 
-template<unsigned BSIZE> void memory_space_impl<BSIZE>::read( mem_addr_t addr, size_t length, void *data ) const
+template<unsigned BSIZE> void memory_space_impl<BSIZE>::read_single_block( mem_addr_t blk_idx, mem_addr_t addr, size_t length, void *data) const
 {
-   mem_addr_t index = addr >> m_log2_block_size;
-   unsigned offset = addr & (BSIZE-1);
-   unsigned nbytes = length;
-   typename map_t::const_iterator i = m_data.find(index);
-   if( (addr+length) > (index+1)*BSIZE ) {
+   if ((addr + length) > (blk_idx + 1) * BSIZE) {
       printf("GPGPU-Sim PTX: ERROR * access to memory \'%s\' is unaligned : addr=0x%x, length=%zu\n",
              m_name.c_str(), addr, length);
       printf("GPGPU-Sim PTX: (addr+length)=0x%lx > 0x%x=(index+1)*BSIZE, index=0x%x, BSIZE=0x%x\n",
-             (addr+length),(index+1)*BSIZE, index, BSIZE);
+             (addr+length),(blk_idx+1)*BSIZE, blk_idx, BSIZE);
       throw 1;
    }
+   typename map_t::const_iterator i = m_data.find(blk_idx);
    if( i == m_data.end() ) {
       for( size_t n=0; n < length; n++ ) 
          ((unsigned char*)data)[n] = (unsigned char) 0;
       //printf("GPGPU-Sim PTX:  WARNING reading %zu bytes from unititialized memory at address 0x%x in space %s\n", length, addr, m_name.c_str() );
    } else {
+      unsigned offset = addr & (BSIZE-1);
+      unsigned nbytes = length;
       i->second.read(offset,nbytes,(unsigned char*)data);
    }
 }
 
+template<unsigned BSIZE> void memory_space_impl<BSIZE>::read( mem_addr_t addr, size_t length, void *data ) const
+{
+   mem_addr_t index = addr >> m_log2_block_size;
+   if ((addr+length) <= (index+1)*BSIZE ) {
+      // fast route for intra-block access 
+      read_single_block(index, addr, length, data); 
+   } else {
+      // slow route for inter-block access 
+      unsigned nbytes_remain = length;
+      unsigned dst_offset = 0; 
+      mem_addr_t current_addr = addr; 
+
+      while (nbytes_remain > 0) {
+         unsigned offset = current_addr & (BSIZE-1);
+         mem_addr_t page = current_addr >> m_log2_block_size; 
+         mem_addr_t access_limit = offset + nbytes_remain; 
+         if (access_limit > BSIZE) {
+            access_limit = BSIZE;
+         } 
+         
+         size_t tx_bytes = access_limit - offset; 
+         read_single_block(page, current_addr, tx_bytes, &((unsigned char*)data)[dst_offset]); 
+
+         // advance pointers 
+         dst_offset += tx_bytes; 
+         current_addr += tx_bytes; 
+         nbytes_remain -= tx_bytes; 
+      }
+      assert(nbytes_remain == 0); 
+   }
+}
+
 template<unsigned BSIZE> void memory_space_impl<BSIZE>::print( const char *format, FILE *fout ) const
 {
    typename map_t::const_iterator i_page;
diff --git a/src/cuda-sim/memory.h b/src/cuda-sim/memory.h
index bac4d24..4487b36 100644
--- a/src/cuda-sim/memory.h
+++ b/src/cuda-sim/memory.h
@@ -164,6 +164,7 @@ public:
    virtual void set_watch( addr_t addr, unsigned watchpoint ); 
 
 private:
+   void read_single_block( mem_addr_t blk_idx, mem_addr_t addr, size_t length, void *data) const; 
    std::string m_name;
    unsigned m_log2_block_size;
    typedef mem_map<mem_addr_t,mem_storage<BSIZE> > map_t;
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index c1217ee..5f9caaa 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -77,7 +77,7 @@ 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(ATOM_OP,atom_impl,"atom",1,3)
 OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",1,3)
 OP_DEF(BFE_OP,bfe_impl,"bfe",1,1)
 OP_DEF(BFI_OP,bfi_impl,"bfi",1,1)
@@ -86,6 +86,7 @@ OP_DEF(BRA_OP,bra_impl,"bra",0,3)
 OP_DEF(BREV_OP,brev_impl,"brev",1,1)
 OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9)
 OP_DEF(CALL_OP,call_impl,"call",1,3)
+OP_DEF(CALLP_OP,callp_impl,"callp",1,3)
 OP_DEF(CLZ_OP,clz_impl,"clz",1,1)
 OP_DEF(CNOT_OP,cnot_impl,"cnot",1,1)
 OP_DEF(COS_OP,cos_impl,"cos",1,4)
@@ -122,6 +123,7 @@ OP_DEF(RCP_OP,rcp_impl,"rcp",1,4)
 OP_DEF(RED_OP,red_impl,"red",1,7)
 OP_DEF(REM_OP,rem_impl,"rem",1,1)
 OP_DEF(RET_OP,ret_impl,"ret",0,3)
+OP_DEF(RETP_OP,retp_impl,"retp",0,3)
 OP_DEF(RSQRT_OP,rsqrt_impl,"rsqrt",1,4)
 OP_DEF(SAD_OP,sad_impl,"sad",1,1)
 OP_DEF(SELP_OP,selp_impl,"selp",1,1)
diff --git a/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc
index 567c62b..40e574a 100644
--- a/src/cuda-sim/ptx-stats.cc
+++ b/src/cuda-sim/ptx-stats.cc
@@ -67,9 +67,6 @@
 #include <map>
 #include "../tr1_hash_map.h"
 
-// external dependencies
-extern function_info *g_func_info;
-
 // options
 bool enable_ptx_file_line_stats;
 char * ptx_line_stats_filename = NULL;
@@ -192,7 +189,7 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn)
 
 // attribute pipeline latency to this ptx instruction (specified by the pc)
 // pipeline latency is the number of cycles a warp with this instruction spent in the pipeline
-void ptx_file_line_stats_add_latency(void * ptx_thd, unsigned pc, unsigned latency)
+void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency)
 {
     const ptx_instruction *pInsn = function_info::pc_to_instruction(pc);
     
@@ -210,7 +207,7 @@ void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic)
 
 // attribute the number of shared memory access cycles to a ptx instruction
 // counts both the number of warps doing shared memory access and the number of cycles involved
-void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, unsigned n_way_bkconflict)
+void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict)
 {
     const ptx_instruction *pInsn = function_info::pc_to_instruction(pc);
     
@@ -221,7 +218,7 @@ void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, uns
 
 // attribute a non-coalesced mem access to a ptx instruction 
 // counts both the number of warps causing this and the number of memory requests generated
-void ptx_file_line_stats_add_uncoalesced_gmem(void * ptx_thd, unsigned pc, unsigned n_access)
+void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access)
 {
     const ptx_instruction *pInsn = function_info::pc_to_instruction(pc);
     
diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h
index 004433b..019586c 100644
--- a/src/cuda-sim/ptx-stats.h
+++ b/src/cuda-sim/ptx-stats.h
@@ -78,10 +78,10 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn);
 #endif
 
 // stat collection interface to gpgpu-sim
-void ptx_file_line_stats_add_latency(void * ptx_thd, unsigned pc, unsigned latency);
+void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency);
 void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic);
-void ptx_file_line_stats_add_smem_bank_conflict(void * ptx_thd, unsigned pc, unsigned n_way_bkconflict);
-void ptx_file_line_stats_add_uncoalesced_gmem(void * ptx_thd, unsigned pc, unsigned n_access);
+void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict);
+void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access);
 
 void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores);
 void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc);
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 82daef9..1b87c7d 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -98,6 +98,7 @@ bra     TC; ptx_lval.int_value = BRA_OP; return OPCODE;
 brev    TC; ptx_lval.int_value = BREV_OP; return OPCODE;
 brkpt   TC; ptx_lval.int_value = BRKPT_OP; return OPCODE;
 call	TC; ptx_lval.int_value = CALL_OP; return OPCODE;
+callp    TC; ptx_lval.int_value = CALLP_OP; return OPCODE;
 clz	TC; ptx_lval.int_value = CLZ_OP; return OPCODE;
 cnot	TC; ptx_lval.int_value = CNOT_OP; return OPCODE;
 cos	TC; ptx_lval.int_value = COS_OP; return OPCODE;
@@ -135,6 +136,7 @@ rcp	TC; ptx_lval.int_value = RCP_OP; return OPCODE;
 red     TC; ptx_lval.int_value = RED_OP; return OPCODE;
 rem	TC; ptx_lval.int_value = REM_OP; return OPCODE;
 ret	TC; ptx_lval.int_value = RET_OP; return OPCODE;
+retp     TC; ptx_lval.int_value = RETP_OP; return OPCODE;
 rsqrt	TC; ptx_lval.int_value = RSQRT_OP; return OPCODE;
 sad     TC; ptx_lval.int_value = SAD_OP; return OPCODE;
 selp	TC; ptx_lval.int_value = SELP_OP; return OPCODE;
@@ -263,7 +265,7 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.bb128  TC;  return BB128_TYPE;
 \.pred TC;  return PRED_TYPE;
 
-\.texref TC;  return TEXREF_TYPE;
+\.texref TC; BEGIN(NOT_OPCODE); return TEXREF_TYPE;
 \.samplerref TC;  return SAMPLERREF_TYPE;
 \.surfref TC;  return SURFREF_TYPE;
 
@@ -271,6 +273,8 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.v3	TC; return V3_TYPE;
 \.v4	TC; return V4_TYPE;
 
+\.half  TC; return HALF_OPTION; /* ptxplus */
+
 \.equ	TC; return EQU_OPTION;
 \.neu	TC; return NEU_OPTION;
 \.ltu	TC; return LTU_OPTION;
@@ -322,6 +326,7 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.all	TC; return ALL_OPTION;
 \.gl    TC; return GLOBAL_OPTION;
 \.cta   TC; return CTA_OPTION;
+\.sys   TC; return SYS_OPTION;
 
 \.exit   TC; return EXIT_OPTION;
 
@@ -368,11 +373,11 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 ">"	TC; return RIGHT_ANGLE_BRACKET;
 "("	TC; return LEFT_PAREN;
 ")"	TC; return RIGHT_PAREN;
-":"	TC; return COLON;
+":"	TC; BEGIN(INITIAL); return COLON;
 ";"	TC; BEGIN(INITIAL); return SEMI_COLON;
 "!"	TC; return EXCLAMATION;
 "=" 	TC; return EQUALS;
-"{"	TC; BEGIN(INITIAL); return LEFT_BRACE;
+"{"	TC; return LEFT_BRACE;
 "}"	TC; return RIGHT_BRACE;
 \.	TC; return PERIOD;
 "/"	TC; return BACKSLASH;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 1551f22..833e082 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -133,6 +133,7 @@
 %token  V4_TYPE
 %token  COMMA
 %token  PRED
+%token  HALF_OPTION
 %token  EQ_OPTION
 %token  NE_OPTION
 %token  LT_OPTION
@@ -205,6 +206,7 @@
 %token  ALL_OPTION
 %token  GLOBAL_OPTION
 %token  CTA_OPTION
+%token  SYS_OPTION
 %token  EXIT_OPTION
 %token  ABS_OPTION
 %token  TO_OPTION
@@ -260,7 +262,8 @@ function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; func_header(".ent
 	| EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; func_header(".func"); }
 	;
 
-param_list: param_entry { add_directive(); }
+param_list: /*empty*/
+	| 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(); }
@@ -277,8 +280,10 @@ statement_list: directive_statement { add_directive(); }
 	;
 
 directive_statement: variable_declaration SEMI_COLON
-	| VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2); }
+	| VERSION_DIRECTIVE DOUBLE_OPERAND { add_version_info($2, 0); }
+	| VERSION_DIRECTIVE DOUBLE_OPERAND PLUS { add_version_info($2,1); }
 	| TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER { target_header2($2,$4); }
+	| TARGET_DIRECTIVE IDENTIFIER COMMA IDENTIFIER COMMA IDENTIFIER { target_header3($2,$4,$6); }
 	| TARGET_DIRECTIVE IDENTIFIER { target_header($2); }
 	| FILE_DIRECTIVE INT_OPERAND STRING { add_file($2,$3); } 
 	| LOC_DIRECTIVE INT_OPERAND INT_OPERAND INT_OPERAND 
@@ -416,6 +421,7 @@ option: type_spec
 	| ALL_OPTION { add_option(ALL_OPTION); }
 	| GLOBAL_OPTION { add_option(GLOBAL_OPTION); }
 	| CTA_OPTION { add_option(CTA_OPTION); }
+	| SYS_OPTION { add_option(SYS_OPTION); }
 	| GEOM_MODIFIER_1D { add_option(GEOM_MODIFIER_1D); }
 	| GEOM_MODIFIER_2D { add_option(GEOM_MODIFIER_2D); }
 	| GEOM_MODIFIER_3D { add_option(GEOM_MODIFIER_3D); }
@@ -428,6 +434,7 @@ option: type_spec
 	| ABS_OPTION { add_option(ABS_OPTION); }
 	| atomic_operation_spec ;
 	| TO_OPTION { add_option(TO_OPTION); }
+	| HALF_OPTION { add_option(HALF_OPTION); }
 	;
 
 atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } 
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index cd925a5..db915fe 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -118,10 +118,10 @@ symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol
 {
    m_scope_name = std::string(scope_name);
    m_reg_allocator=0;
-   m_shared_next = 0x100; // for debug with valgrind: make zero imply undefined address
-   m_const_next  = 0x100; // for debug with valgrind: make zero imply undefined address
-   m_global_next = 0x100; // for debug with valgrind: make zero imply undefined address
-   m_local_next  = 0x100; // for debug with valgrind: make zero imply undefined address
+   m_shared_next = 0x100;
+   m_const_next  = 0x100;
+   m_global_next = 0x100;
+   m_local_next  = 0;
    m_parent = parent;
    if ( m_parent ) {
       m_shared_next = m_parent->m_shared_next;
@@ -140,12 +140,24 @@ const ptx_version &symbol_table::get_ptx_version() const
    else return m_parent->get_ptx_version(); 
 }
 
+unsigned symbol_table::get_sm_target() const 
+{ 
+   if( m_parent == NULL ) 
+      return m_ptx_version.target();
+   else return m_parent->get_sm_target(); 
+}
+
 void symbol_table::set_ptx_version( float ver, unsigned ext ) 
 { 
    m_ptx_version = ptx_version(ver,ext); 
    assert( m_ptx_version.ver() < 3 );
 }
 
+void symbol_table::set_sm_target( const char *target, const char *ext, const char *ext2 )
+{
+   m_ptx_version.set_target(target,ext,ext2);
+}
+
 symbol *symbol_table::lookup( const char *identifier ) 
 {
    std::string key(identifier);
@@ -306,13 +318,13 @@ void function_info::create_basic_blocks()
          i = find_next_real_instruction(++i);
       } else {
          switch( pI->get_opcode() ) {
-         case BRA_OP: case RET_OP: case EXIT_OP:
+         case BRA_OP: case RET_OP: case EXIT_OP: case RETP_OP:
             i++;
             if( i != m_instructions.end() ) 
                leaders.push_back(*i);
             i = find_next_real_instruction(i);
             break;
-         case CALL_OP:
+         case CALL_OP: case CALLP_OP:
             if( pI->has_pred() ) {
                printf("GPGPU-Sim PTX: Warning found predicated call\n");
                i++;
@@ -424,7 +436,7 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f
       ptx_instruction *pI = (*bb_itr)->ptx_end;
       if ((*bb_itr)->is_exit) //reached last basic block, no successors to link 
          continue;
-      if (pI->get_opcode() == RET_OP || pI->get_opcode() == EXIT_OP ) {
+      if (pI->get_opcode() == RETP_OP || pI->get_opcode() == RET_OP || pI->get_opcode() == EXIT_OP ) {
          (*bb_itr)->successor_ids.insert(exit_bb->bb_id);
          exit_bb->predecessor_ids.insert((*bb_itr)->bb_id);
          if( pI->has_pred() ) {
@@ -451,7 +463,7 @@ void function_info::connect_basic_blocks( ) //iterate across m_basic_blocks of f
          // if basic block does not end in an unpredicated branch, 
          // then next basic block is also successor
          // (this is better than testing for .uni)
-         unsigned next_addr = pI->get_m_instr_mem_index() + 1;
+         unsigned next_addr = pI->get_m_instr_mem_index() + pI->inst_size();
          basic_block_t *next_bb = m_instr_mem[next_addr]->get_bb();
          (*bb_itr)->successor_ids.insert(next_bb->bb_id);
          next_bb->predecessor_ids.insert((*bb_itr)->bb_id);
@@ -653,12 +665,12 @@ void function_info::find_ipostdominators( )
 
 void function_info::find_idominators( )
 {  
-   // find immediate postdominator blocks, using algorithm of
+   // find immediate dominator 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() );
+   printf("GPGPU-Sim PTX: Finding immediate dominators 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
+   assert( m_basic_blocks.size() >= 2 ); // must have a distinquished entry block
+   for (unsigned i=0; i<m_basic_blocks.size(); i++) { //initialize Tmp(n) to all doms 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);
@@ -683,9 +695,9 @@ void function_info::find_idominators( )
    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
+      //assert( m_basic_blocks[n]->Tmp_ids.size() <= 1 );
+         // if the above assert fails we have an error in either dominator
+         // computation, the flow graph does not have a unique entry, 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++;
@@ -843,6 +855,8 @@ unsigned type_info_key::type_decode( int type, size_t &size, int &basic_type )
    case B16_TYPE: size=16; basic_type=0; return 13;
    case B32_TYPE: size=32; basic_type=0; return 14;
    case B64_TYPE: size=64; basic_type=0; return 15;
+   case BB64_TYPE: size=64; basic_type=0; return 15;
+   case BB128_TYPE: size=128; basic_type=0; return 16;
    case TEXREF_TYPE: case SAMPLERREF_TYPE: case SURFREF_TYPE:
       size=32; basic_type=3; return 16;
    default: 
@@ -968,7 +982,8 @@ ptx_instruction::ptx_instruction( int opcode,
                                   memory_space_t space_spec,
                                   const char *file, 
                                   unsigned line,
-                                  const char *source ) 
+                                  const char *source,
+                                  unsigned warp_size ) : inst_t()
 {
    m_uid = ++g_num_ptx_inst_uid;
    m_PC = 0;
@@ -996,8 +1011,9 @@ ptx_instruction::ptx_instruction( int opcode,
    m_geom_spec = 0;
    m_vector_spec = 0;
    m_atomic_spec = 0;
-   m_warp_size = ::warp_size;
+   m_warp_size = warp_size;
    m_membar_level = 0;
+   m_inst_size = 8; // bytes
 
    std::list<int>::const_iterator i;
    unsigned n=1;
@@ -1096,6 +1112,9 @@ ptx_instruction::ptx_instruction( int opcode,
       case CTA_OPTION:
          m_membar_level = CTA_OPTION;
          break;
+      case SYS_OPTION:
+         m_membar_level = SYS_OPTION;
+         break;
       case FTZ_OPTION:
          break;
       case EXIT_OPTION:
@@ -1113,7 +1132,9 @@ ptx_instruction::ptx_instruction( int opcode,
       case CA_OPTION: case CG_OPTION: case CS_OPTION: case LU_OPTION: case CV_OPTION:
          m_cache_option = last_ptx_inst_option;
          break;
-
+      case HALF_OPTION:
+         m_inst_size = 4; // bytes
+         break;
       default:
          assert(0);
          break;
@@ -1141,7 +1162,7 @@ void ptx_instruction::print_insn( FILE *fp ) const
    snprintf(buf,1024,"%s", m_source.c_str());
    p = strtok(buf,";");
    if( !is_label() ) 
-      fprintf(fp,"PC=%3u [idx=%3u] ", m_PC, m_instr_mem_index );
+      fprintf(fp," PC=%3u [%3u] ", m_PC, m_instr_mem_index );
    else
       fprintf(fp,"                " );
    fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p );
@@ -1168,7 +1189,13 @@ function_info::function_info(int entry_point )
 void function_info::print_insn( unsigned pc, FILE * fp ) const
 {
    unsigned index = pc - m_start_PC;
-   fprintf(fp,"FUNC[%s]",m_name.c_str() );
+   char command[1024];
+   char buffer[1024];
+   snprintf(command,1024,"c++filt -p %s",m_name.c_str());
+   FILE *p = popen(command,"r");
+   buffer[0]=0;
+   fscanf(p,"%1023s",buffer);
+   fprintf(fp,"%s",buffer);
    if ( index >= m_instr_mem_size ) {
       fprintf(fp, "<past last instruction (max pc=%u)>", m_start_PC + m_instr_mem_size - 1 );
    } else {
@@ -1191,10 +1218,13 @@ extern "C" FILE *ptxinfo_in;
 void gpgpu_ptx_assemble( std::string kname, void *kinfo )
 {
     function_info *func_info = (function_info *)kinfo;
+    if((function_info *)kinfo == NULL) {
+       printf("GPGPU-Sim PTX: Error - missing function definition \'%s\'\n", kname.c_str());
+       abort();
+    }
     if( func_info->is_extern() ) {
        printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() );
        return;
     }
-    g_func_info = func_info;
     func_info->ptx_assemble();
 }
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 43ac92a..afb2476 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -319,7 +319,9 @@ public:
    symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent );
    void set_name( const char *name );
    const ptx_version &get_ptx_version() const;
+   unsigned get_sm_target() const;
    void set_ptx_version( float ver, unsigned ext );
+   void set_sm_target( const char *target, const char *ext, const char *ext2 );
    symbol* lookup( const char *identifier );
    std::string get_scope_name() const { return m_scope_name; }
    symbol *add_variable( const char *identifier, const type_info *type, unsigned size, const char *filename, unsigned line );
@@ -784,9 +786,9 @@ struct gpgpu_recon_t {
    address_type target_pc;
 };
 
-class ptx_instruction {
+class ptx_instruction : public inst_t {
 public:
-   ptx_instruction( int opcode, 
+    ptx_instruction( int opcode, 
                     const symbol *pred, 
                     int neg_pred, 
                     int pred_mod,
@@ -798,9 +800,13 @@ public:
                     memory_space_t space_spec,
                     const char *file, 
                     unsigned line,
-                    const char *source );
+                    const char *source,
+                    unsigned warp_size );
+
+
    void print_insn() const;
-   void print_insn( FILE *fp ) const;
+   virtual void print_insn( FILE *fp ) const;
+   unsigned inst_size() const { return m_inst_size; }
    unsigned uid() const { return m_uid;}
    int get_opcode() const { return m_opcode;}
    const char *get_opcode_cstr() const 
@@ -961,6 +967,7 @@ public:
    }
 
 private:
+
    basic_block_t        *m_basic_block;
    unsigned          m_uid;
    addr_t            m_PC;
@@ -999,6 +1006,10 @@ private:
    enum vote_mode_t m_vote_mode;
    int m_membar_level;
    int m_instr_mem_index; //index into m_instr_mem array
+   unsigned m_inst_size; // bytes
+
+   virtual void pre_decode();
+   friend class function_info;
 };
 
 class param_info {
@@ -1036,6 +1047,7 @@ class function_info {
 public:
    function_info(int entry_point );
    const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); }
+   unsigned get_sm_target() const { return m_symtab->get_sm_target(); }
    bool is_extern() const { return m_extern; }
    void set_name(const char *name)
    {
@@ -1094,22 +1106,7 @@ public:
    unsigned get_function_size() { return m_instructions.size();}
 
    void ptx_assemble();
-   void ptx_decode_inst( ptx_thread_info *thd, 
-                         unsigned *op_type, 
-                         int *i1, 
-                         int *i2, 
-                         int *i3,
-                         int *i4,
-                         int *o1,
-                         int *o2,
-                         int *o3,
-                         int *o4,
-                         int *vectorin,
-                         int *vectorout,
-                         int *arch_reg,
-                         int *pred,
-                         int *ar1, int *ar2 );
-   void ptx_exec_inst( ptx_thread_info *thd, addr_t *addr, memory_space_t *space, unsigned *data_size, unsigned *cycles, dram_callback_t* callback, unsigned warp_active_mask  );
+ 
    unsigned ptx_get_inst_op( ptx_thread_info *thread );
    void add_param( const char *name, struct param_t value )
    {
@@ -1169,6 +1166,8 @@ public:
 
    const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info &info) {
       m_kernel_info = info;
+      m_kernel_info.ptx_version = 10*get_ptx_version().ver();
+      m_kernel_info.sm_target = get_ptx_version().target();
    }
    symbol_table *get_symtab()
    {
@@ -1177,9 +1176,8 @@ public:
 
    static const ptx_instruction* pc_to_instruction(unsigned pc) 
    {
-      assert(pc > 0);
       assert(pc <= s_g_pc_to_insn.size());
-      return s_g_pc_to_insn[pc - 1];
+      return s_g_pc_to_insn[pc];
    }
    unsigned local_mem_framesize() const 
    { 
@@ -1404,16 +1402,12 @@ struct textureInfo {
    unsigned int texel_size_numbits; //log2(texel_size)
 };
 
-
-extern function_info *g_func_info;
-
-extern function_info *g_entrypoint_func_info;
 extern std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 
-#define GLOBAL_HEAP_START 0x10000000
+#define GLOBAL_HEAP_START 0x80000000
    // start allocating from this address (lower values used for allocating globals in .ptx file)
 #define SHARED_MEM_SIZE_MAX (64*1024)
-#define LOCAL_MEM_SIZE_MAX 1024
+#define LOCAL_MEM_SIZE_MAX (16*1024)
 #define MAX_STREAMING_MULTIPROCESSORS 64
 #define MAX_THREAD_PER_SM 1024
 #define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index 783ce91..1aa5101 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -67,6 +67,7 @@
 #include "cuda-sim.h"
 #include "ptx_parser.h"
 #include <dirent.h>
+#include <fstream>
 
 /// globals
 
@@ -78,6 +79,7 @@ bool g_override_embedded_ptx = false;
 
 struct ptx_info_t {
     char *str;
+    char *cubin_str;
     char *fname;
     ptx_info_t *next;
 };
@@ -93,6 +95,9 @@ extern "C" int ptxinfo_parse();
 extern "C" int ptxinfo_debug;
 extern "C" FILE *ptxinfo_in;
 
+extern int g_ptx_convert_to_ptxplus;
+extern int g_ptx_save_converted_ptxplus;
+
 /// static functions
 
 static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr) 
@@ -223,11 +228,23 @@ void gpgpu_ptx_sim_load_gpu_kernels()
             printf("GPGPU-Sim PTX: USING EMBEDDED .ptx files...\n"); 
             ptx_info_t *s;
             for ( s=g_ptx_source_array; s!=NULL; s=s->next ) {
-                 symbol_table *symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num);
+            	 symbol_table *symtab;
+            	 source_num++;
+            	 if(g_ptx_convert_to_ptxplus) {
+            		 char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(s->str, s->cubin_str, source_num);
+                     symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, s->str, source_num);
+                     delete[] ptxplus_str;
+            	 } else {
+                     symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, s->str, source_num);
+            	 }
                  load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
                  load_constants(symtab,STATIC_ALLOC_LIMIT);
             }
         } else {
+        	if(g_ptx_convert_to_ptxplus) {
+        		perror("GPGPU-Sim PTX: convert_to_ptxplus option enabled. Cannot use this option with external ptx files.\n");
+        		assert(0);
+        	}
             const char *filename = NULL;
             struct dirent **namelist;
             int n = scandir(".", &namelist, ptx_file_filter, alphasort);
@@ -235,20 +252,17 @@ void gpgpu_ptx_sim_load_gpu_kernels()
                 perror("GPGPU-Sim PTX: no PTX files returned by scandir");
             else {
                 while (n--) {
-                    if ( filename != NULL ) {
-                        printf("Loader error: support for multiple .ptx files not yet enabled\n");
-                        abort();
-                    }
                     filename = namelist[n]->d_name;
                     printf("Parsing %s..\n", filename);
                     ptx_in = fopen( filename, "r" );
-                    free(namelist[n]);
                     symbol_table *symtab=init_parser(filename);
                     ptx_parse ();
                     ptxinfo_in = open_ptxinfo(filename);
                     ptxinfo_parse();
                     load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF);
                     load_constants(symtab,STATIC_ALLOC_LIMIT);
+
+                    free(namelist[n]);
                 }
                 free(namelist);
             }
@@ -278,11 +292,17 @@ void gpgpu_ptx_sim_load_gpu_kernels()
    }
 }
 
-void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname )
+void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *cubin_string, const char *sourcefname )
 {
     ptx_info_t *t = new ptx_info_t;
     t->next = NULL;
     t->str = strdup(ptx_string);
+    if (cubin_string != NULL) {
+       t->cubin_str = strdup(cubin_string);
+    } else {
+       assert(g_ptx_convert_to_ptxplus == 0); 
+       t->cubin_str = NULL; 
+    }
     t->fname = strdup(sourcefname);
 
     // put ptx source into a fifo
@@ -338,7 +358,106 @@ void print_ptx_file( const char *p, unsigned source_num, const char *filename )
    fflush(stdout);
 }
 
-symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) 
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num)
+{
+	printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
+
+	// Extract ptx to a file
+    char fname_ptx[1024];
+    snprintf(fname_ptx,1024,"_ptx_XXXXXX");
+    int fd=mkstemp(fname_ptx);
+    close(fd);
+
+    printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname_ptx);
+    FILE *ptxfile = fopen(fname_ptx,"w");
+    fprintf(ptxfile,"%s",ptx_str);
+    fclose(ptxfile);
+
+    // Extract cubin to a file
+    char fname_cubin[1024];
+    snprintf(fname_cubin,1024,"_cubin_XXXXXX");
+    int fd2=mkstemp(fname_cubin);
+    close(fd2);
+
+    printf("GPGPU-Sim PTX: extracting embedded cubin to temporary file \"%s\"\n", fname_cubin);
+    FILE *cubinfile = fopen(fname_cubin,"w");
+    fprintf(cubinfile,"%s",cubin_str);
+    fclose(cubinfile);
+
+    // Run decuda
+    char fname_decuda[1024];
+    snprintf(fname_decuda,1024,"_decuda_XXXXXX");
+    int fd3=mkstemp(fname_decuda);
+    close(fd3);
+
+    char decuda_commandline[1024];
+    snprintf(decuda_commandline,1024,"$DECUDA_INSTALL_PATH/decuda.py -o %s %s", fname_decuda, fname_cubin);
+
+    printf("GPGPU-Sim PTX: calling decuda on cubin file, decuda output file = \"%s\"\n", fname_decuda);
+    int decuda_result = system(decuda_commandline);
+	if( decuda_result != 0 ) {
+	   printf("GPGPU-Sim PTX: ERROR ** while calling decuda (b) %d\n", decuda_result);
+	   printf("               Ensure env variable DECUDA_INSTALL_PATH is set and points to decuda base directory.\n");
+	   exit(1);
+	}
+
+	// Run decuda_to_ptxplus
+    char fname_ptxplus[1024];
+    snprintf(fname_ptxplus,1024,"_ptxplus_XXXXXX");
+    int fd4=mkstemp(fname_ptxplus);
+    close(fd4);
+
+    char d2pp_commandline[1024];
+    snprintf(d2pp_commandline,1024,"$D2PP_INSTALL_PATH/decuda_to_ptxplus %s %s %s %s > /dev/null", fname_decuda, fname_ptx, fname_cubin, fname_ptxplus);
+
+    printf("GPGPU-Sim PTX: calling decuda_to_ptxplus, ptxplus output file = \"%s\"\n", fname_ptxplus);
+    int d2pp_result = system(d2pp_commandline);
+	if( d2pp_result != 0 ) {
+	   printf("GPGPU-Sim PTX: ERROR ** while calling decuda_to_ptxplus %d\n", d2pp_result);
+	   printf("               Ensure env variable D2PP_INSTALL_PATH is set and points to decuda_to_ptxplus base directory.\n");
+	   exit(1);
+	}
+
+	// Get ptxplus from file
+	std::ifstream fileStream(fname_ptxplus, std::ios::in);
+	std::string text, line;
+	while(getline(fileStream,line)) {
+		text += (line + "\n");
+	}
+	fileStream.close();
+
+	char* ptxplus_str = new char [strlen(text.c_str())+1];
+	strcpy(ptxplus_str, text.c_str());
+
+	// Save ptxplus to file if specified
+	if(g_ptx_save_converted_ptxplus) {
+	    char fname_ptxplus_save[1024];
+	    snprintf(fname_ptxplus_save,1024,"_%u.ptxplus", source_num );
+		printf("GPGPU-Sim PTX: saving converted ptxplus to file \"%s\"\n", fname_ptxplus_save);
+
+	    FILE *file_ptxplus_save = fopen(fname_ptxplus_save,"w");
+	    fprintf(file_ptxplus_save,"%s",ptxplus_str);
+	    fclose(file_ptxplus_save);
+	}
+
+	// Remove temporary files
+	char rm_commandline[1024];
+	snprintf(rm_commandline,1024,"rm -f %s %s %s %s", fname_ptx, fname_cubin, fname_decuda, fname_ptxplus);
+	printf("GPGPU-Sim PTX: removing temporary files using \"%s\"\n", rm_commandline);
+	int rm_result = system(rm_commandline);
+	if( rm_result != 0 ) {
+	   printf("GPGPU-Sim PTX: ERROR ** while removing temporary files %d\n", rm_result);
+	   exit(1);
+	}
+
+	printf("GPGPU-Sim PTX: DONE converting EMBEDDED .ptx file to ptxplus \n");
+
+	return ptxplus_str;
+
+}
+
+
+symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info,  unsigned source_num )
 {
     char buf[1024];
     snprintf(buf,1024,"_%u.ptx", source_num );
@@ -375,7 +494,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source
 
     printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname);
     FILE *ptxfile = fopen(fname,"w");
-    fprintf(ptxfile,"%s",p);
+    fprintf(ptxfile,"%s", p_for_info);
     fclose(ptxfile);
 
     char fname2[1024];
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index 95130f1..01dc0b7 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -75,7 +75,8 @@ extern memory_space *g_param_mem;
 extern bool g_override_embedded_ptx;
  
 void gpgpu_ptx_sim_load_gpu_kernels();
-void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname );
-class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
+void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *cubin_string,  const char *sourcefname );
+class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num );
+char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num);
 
 #endif
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index d4a892f..5a65481 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -68,6 +68,12 @@
 extern "C" int ptx_error( const char *s );
 extern int ptx_lineno;
 
+static unsigned g_warp_size;
+void set_ptx_warp_size(unsigned warp_size)
+{
+   g_warp_size=warp_size;
+}
+
 static bool g_debug_ir_generation=false;
 const char *g_filename;
 unsigned g_max_regs_per_thread = 0;
@@ -110,12 +116,10 @@ std::list<int> g_scalar_type;
       fflush(stdout); \
    }
 
-unsigned g_entry_func_param_index=0;
-function_info *g_func_info = NULL;
-function_info *g_entrypoint_func_info = NULL;
-symbol_table *g_entrypoint_symbol_table = NULL;
-std::map<unsigned,std::string> g_ptx_token_decode;
-operand_info g_return_var;
+static unsigned g_entry_func_param_index=0;
+static function_info *g_func_info = NULL;
+static std::map<unsigned,std::string> g_ptx_token_decode;
+static operand_info g_return_var;
 
 const char *decode_token( int type )
 {
@@ -194,9 +198,6 @@ void add_function_name( const char *name )
 {
    DPRINTF("add_function_name %s %s", name,  ((g_entry_point==1)?"(entrypoint)":((g_entry_point==2)?"(extern)":"")));
    bool prior_decl = g_global_symbol_table->add_function_decl( name, g_entry_point, &g_func_info, &g_current_symbol_table );
-   if( g_entry_point ) {
-      g_entrypoint_func_info = g_func_info;
-   }
    if( g_add_identifier_cached__identifier ) {
       add_identifier( g_add_identifier_cached__identifier,
                       g_add_identifier_cached__array_dim,
@@ -270,7 +271,7 @@ extern "C" char linebuf[1024];
 
 void set_return()
 {
-   parse_assert( (g_opcode == CALL_OP), "only call can have return value");
+   parse_assert( (g_opcode == CALL_OP || g_opcode == CALLP_OP), "only call can have return value");
    g_operands.front().set_return();
    g_return_var = g_operands.front();
 }
@@ -291,6 +292,7 @@ const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned li
 void add_instruction() 
 {
    DPRINTF("add_instruction: %s", ((g_opcode>0)?g_opcode_string[g_opcode]:"<label>") );
+   assert( g_warp_size != 0 );
    ptx_instruction *i = new ptx_instruction( g_opcode, 
                                              g_pred, 
                                              g_neg_pred,
@@ -303,7 +305,8 @@ void add_instruction()
                                              g_space_spec,
                                              g_filename,
                                              ptx_lineno,
-                                             linebuf );
+                                             linebuf,
+                                             g_warp_size );
    g_instructions.push_back(i);
    g_inst_lookup[g_filename][ptx_lineno] = i;
    init_instruction_state();
@@ -457,7 +460,14 @@ void add_identifier( const char *identifier, int array_dim, unsigned array_ident
       break;
    case local_space:
       if( g_func_info == NULL ) {
-         printf("GPGPU-Sim PTX: not allocating .local \"%s\" declared at global scope\n", identifier);
+	      printf("GPGPU-Sim PTX: allocating local region for \"%s\" from 0x%x to 0x%lx (local memory space)\n",
+             identifier,
+             g_current_symbol_table->get_local_next(),
+             g_current_symbol_table->get_local_next() + num_bits/8 );
+         fflush(stdout);
+         assert( (num_bits%8) == 0  );
+         g_last_symbol->set_address( g_current_symbol_table->get_local_next() );
+         g_current_symbol_table->alloc_local( num_bits/8 );
          break;
       }
       printf("GPGPU-Sim PTX: allocating stack frame region for .local \"%s\" from 0x%x to 0x%lx\n",
@@ -647,6 +657,10 @@ void add_4vector_operand( const char *d1, const char *d2, const char *d3, const
    const symbol *s3 = g_current_symbol_table->lookup(d3);
    const symbol *s4 = g_current_symbol_table->lookup(d4);
    parse_assert( s1 != NULL && s2 != NULL && s3 != NULL && s4 != NULL, "v4 component(s) missing declarations.");
+   const symbol *null_op = g_current_symbol_table->lookup("_");
+   if ( s2 == null_op ) s2 = NULL;
+   if ( s3 == null_op ) s3 = NULL;
+   if ( s4 == null_op ) s4 = NULL;
    g_operands.push_back( operand_info(s1,s2,s3,s4) );
 }
 
@@ -789,7 +803,7 @@ void add_scalar_operand( const char *identifier )
    DPRINTF("add_scalar_operand");
    const symbol *s = g_current_symbol_table->lookup(identifier);
    if ( s == NULL ) {
-      if ( g_opcode == BRA_OP ) {
+      if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) {
          // forward branch target...
          s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_lineno);
       } else {
@@ -828,9 +842,9 @@ void add_array_initializer()
    g_last_symbol->add_initializer(g_operands);
 }
 
-void add_version_info( float ver )
+void add_version_info( float ver, unsigned ext )
 {
-   g_global_symbol_table->set_ptx_version(ver,0);
+   g_global_symbol_table->set_ptx_version(ver,ext);
 }
 
 void add_file( unsigned num, const char *filename )
@@ -878,8 +892,21 @@ void add_pragma( const char *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 target_header(char* a) 
+{
+   g_global_symbol_table->set_sm_target(a,NULL,NULL);
+}
+
+void target_header2(char* a, char* b) 
+{
+   g_global_symbol_table->set_sm_target(a,b,NULL);
+}
+
+void target_header3(char* a, char* b, char* c) 
+{
+   g_global_symbol_table->set_sm_target(a,b,c);
+}
 
 void func_header(char* a) {} //intentional dummy function
 void func_header_info(char* a) {} //intentional dummy function
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 4852e39..dba88ca 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -110,7 +110,7 @@ void set_return();
 void add_alignment_spec( int spec );
 void add_array_initializer();
 void add_file( unsigned num, const char *filename );
-void add_version_info( float ver );
+void add_version_info( float ver, unsigned ext);
 void *reset_symtab();
 void set_symtab(void*);
 void add_pragma( const char *str );
@@ -120,6 +120,7 @@ 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 target_header3(char* a, char* b, char* c);
 void add_double_operand( const char *d1, const char *d2 );
 void change_memory_addr_space( const char *identifier );
 void change_operand_lohi( int lohi );
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 6dea2e3..2cd0ba6 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -71,7 +71,6 @@ void feature_not_implemented( const char *f );
 
 std::set<unsigned long long> g_ptx_cta_info_sm_idx_used;
 unsigned long long g_ptx_cta_info_uid = 1;
-extern int gpgpu_option_spread_blocks_across_cores;
 
 ptx_cta_info::ptx_cta_info( unsigned sm_idx )
 {
@@ -276,7 +275,11 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
       return (gpu_sim_cycle + gpu_tot_sim_cycle)*2;
    case CTAID_REG:
       assert( dim_mod < 3 );
-      return m_ctaid[dim_mod];
+      if( dim_mod == 0 ) return m_ctaid.x;
+      if( dim_mod == 1 ) return m_ctaid.y;
+      if( dim_mod == 2 ) return m_ctaid.z;
+      abort();
+      break;
    case ENVREG_REG: feature_not_implemented( "%envreg" ); return 0;
    case GRIDID_REG:
       return m_gridid;
@@ -288,16 +291,28 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
    case LANEMASK_GT_REG: feature_not_implemented( "%lanemask_gt" ); return 0;
    case NCTAID_REG:
       assert( dim_mod < 3 );
-      return m_nctaid[dim_mod];
+      if( dim_mod == 0 ) return m_nctaid.x;
+      if( dim_mod == 1 ) return m_nctaid.y;
+      if( dim_mod == 2 ) return m_nctaid.z;
+      abort();
+      break;
    case NTID_REG:
       assert( dim_mod < 3 );
-      return m_ntid[dim_mod];
+      if( dim_mod == 0 ) return m_ntid.x;
+      if( dim_mod == 1 ) return m_ntid.y;
+      if( dim_mod == 2 ) return m_ntid.z;
+      abort();
+      break;
    case NWARPID_REG: feature_not_implemented( "%nwarpid" ); return 0;
    case PM_REG: feature_not_implemented( "%pm" ); return 0;
    case SMID_REG: feature_not_implemented( "%smid" ); return 0;
    case TID_REG:
       assert( dim_mod < 3 );
-      return m_tid[dim_mod];
+      if( dim_mod == 0 ) return m_tid.x;
+      if( dim_mod == 1 ) return m_tid.y;
+      if( dim_mod == 2 ) return m_tid.z;
+      abort();
+      break;
    case WARPSZ_REG: feature_not_implemented( "WARP_SZ" ); return 0;
    default:
       assert(0);
@@ -312,13 +327,13 @@ 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)
+void ptx_thread_info::cpy_tid_to_reg( dim3 tid )
 {
    //copies %tid.x, %tid.y and %tid.z into $r0
    ptx_reg_t data;
    data.s64=0;
 
-   data.u32=(x + (y<<16) + (z<<26));
+   data.u32=(tid.x + (tid.y<<16) + (tid.z<<26));
 
    const symbol *r0 = m_symbol_table->lookup("$r0");
    set_reg(r0,data);
@@ -386,6 +401,21 @@ void ptx_thread_info::callstack_push( unsigned pc, unsigned rpc, const symbol *r
    m_local_mem_stack_pointer += m_func_info->local_mem_framesize(); 
 }
 
+//ptxplus version of callstack_push.
+void ptx_thread_info::callstack_push_plus( unsigned pc, unsigned rpc, const symbol *return_var_src, const symbol *return_var_dst, unsigned call_uid )
+{
+   m_RPC = -1;
+   m_RPC_updated = true;
+   m_last_was_call = true;
+   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();
+}
+
+
 bool ptx_thread_info::callstack_pop()
 {
    const symbol *rv_src = m_callstack.back().m_return_var_src;
@@ -419,6 +449,40 @@ bool ptx_thread_info::callstack_pop()
    return m_callstack.empty();
 }
 
+//ptxplus version of callstack_pop
+bool ptx_thread_info::callstack_pop_plus()
+{
+   const symbol *rv_src = m_callstack.back().m_return_var_src;
+   const symbol *rv_dst = m_callstack.back().m_return_var_dst;
+   assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value
+
+   // read return value from callee frame
+   arg_buffer_t buffer;
+   if( rv_src != NULL )
+      buffer = copy_arg_to_buffer(this, operand_info(rv_src), rv_dst );
+
+   m_symbol_table = m_callstack.back().m_symbol_table;
+   m_NPC = m_callstack.back().m_PC;
+   m_RPC_updated = true;
+   m_last_was_call = false;
+   m_RPC = m_callstack.back().m_RPC;
+   m_func_info = m_callstack.back().m_func_info;
+   if( m_func_info ) {
+      assert( m_local_mem_stack_pointer >= m_func_info->local_mem_framesize() );
+      m_local_mem_stack_pointer -= m_func_info->local_mem_framesize();
+   }
+   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 )
+      copy_buffer_to_frame(this, buffer);
+
+   return m_callstack.empty();
+}
+
 void ptx_thread_info::dump_callstack() const
 {
    std::list<stack_entry>::const_iterator c=m_callstack.begin();
@@ -476,7 +540,7 @@ void ptx_thread_info::dump_regs( FILE *fp )
       const symbol *sym = r->first;
       ptx_reg_t value = r->second;
       std::string name = sym->name();
-      print_reg(name,value,m_symbol_table);
+      print_reg(fp,name,value,m_symbol_table);
    }
 }
 
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 4741769..3a6d9ad 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -71,23 +71,6 @@
 #include "../abstract_hardware_model.h"
 #include "../tr1_hash_map.h" 
 
-
-struct gpgpu_ptx_sim_arg {
-   const void *m_start;
-   size_t m_nbytes;
-   size_t m_offset;
-   struct gpgpu_ptx_sim_arg *m_next;
-};
-
-//Holds properties of the kernel (Kernel's resource use). These will be zero if 
-//the ptxinfo file is not present.
-struct gpgpu_ptx_sim_kernel_info {
-   int lmem;
-   int smem;
-   int cmem;
-   int regs;
-};
-
 #include <assert.h>
 #include "opcodes.h"
 
@@ -240,18 +223,51 @@ public:
          m_valid = false;
          m_ptx_version = 0;
          m_ptx_extensions = 0;
+         m_sm_version_valid=false;
+         m_texmode_unified=true;
+         m_map_f64_to_f32 = true; 
       }
       ptx_version(float ver, unsigned extensions)
       {
          m_valid = true;
          m_ptx_version = ver;
          m_ptx_extensions = extensions;
+         m_sm_version_valid=false;
+         m_texmode_unified=true;
+      }
+      void set_target( const char *sm_ver, const char *ext, const char *ext2 ) 
+      { 
+         assert( m_valid );
+         m_sm_version_str = sm_ver;
+         check_target_extension(ext); 
+         check_target_extension(ext2); 
+         sscanf(sm_ver,"%u",&m_sm_version);
+         m_sm_version_valid=true; 
       }
-      float ver() const { assert(m_valid); return m_ptx_version; }
+      float    ver() const { assert(m_valid); return m_ptx_version; }
+      unsigned target() const { assert(m_valid&&m_sm_version_valid); return m_sm_version; }
       unsigned extensions() const { assert(m_valid); return m_ptx_extensions; }
 private:
+      void check_target_extension( const char *ext ) 
+      {
+         if( ext ) {
+            if( !strcmp(ext,"texmode_independent") ) 
+               m_texmode_unified=false;
+            else if( !strcmp(ext,"texmode_unified") ) 
+               m_texmode_unified=true;
+            else if( !strcmp(ext,"map_f64_to_f32") ) 
+               m_map_f64_to_f32 = true; 
+            else abort();
+         }
+      }
+
       bool     m_valid;
       float    m_ptx_version;
+      unsigned m_sm_version_valid;
+      std::string m_sm_version_str;
+      bool     m_texmode_unified;
+      bool     m_map_f64_to_f32; 
+      unsigned m_sm_version;
       unsigned m_ptx_extensions;
 };
 
@@ -260,6 +276,9 @@ public:
    ~ptx_thread_info();
    ptx_thread_info();
 
+   void ptx_fetch_inst( inst_t &inst ) const;
+   void ptx_exec_inst( inst_t &inst );
+
    const ptx_version &get_ptx_version() const;
    void set_reg( const symbol *reg, const ptx_reg_t &value );
    ptx_reg_t get_reg( const symbol *reg );
@@ -271,8 +290,7 @@ public:
                                    const ptx_reg_t &data1, 
                                    const ptx_reg_t &data2, 
                                    const ptx_reg_t &data3, 
-                                   const ptx_reg_t &data4, 
-                                   unsigned num_elements );
+                                   const ptx_reg_t &data4 );
 
    function_info *func_info()
    {
@@ -285,22 +303,9 @@ public:
       return m_uid;
    }
 
-   dim3 get_ctaid() const
-   {
-      dim3 r;
-      r.x = m_ctaid[0];
-      r.y = m_ctaid[1];
-      r.z = m_ctaid[2];
-      return r;
-   }
-   dim3 get_tid() const
-   {
-      dim3 r;
-      r.x = m_tid[0];
-      r.y = m_tid[1];
-      r.z = m_tid[2];
-      return r;
-   }
+   dim3 get_ctaid() const { return m_ctaid; }
+   dim3 get_tid() const { return m_tid; }
+   class gpgpu_sim *get_gpu() { return m_core->get_gpu(); }
    unsigned get_hw_tid() const { return m_hw_tid;}
    unsigned get_hw_ctaid() const { return m_hw_ctaid;}
    unsigned get_hw_wid() const { return m_hw_wid;}
@@ -335,46 +340,26 @@ public:
 
    void set_single_thread_single_block()
    {
-      m_ntid[0] = 1;
-      m_ntid[1] = 1;
-      m_ntid[2] = 1;
-      m_ctaid[0] = 0;
-      m_ctaid[1] = 0;
-      m_ctaid[2] = 0;
-      m_tid[0] = 0;
-      m_tid[1] = 0;
-      m_tid[2] = 0;
-      m_nctaid[0] = 1;
-      m_nctaid[1] = 1;
-      m_nctaid[2] = 1;
+      m_ntid.x = 1;
+      m_ntid.y = 1;
+      m_ntid.z = 1;
+      m_ctaid.x = 0;
+      m_ctaid.y = 0;
+      m_ctaid.z = 0;
+      m_tid.x = 0;
+      m_tid.y = 0;
+      m_tid.z = 0;
+      m_nctaid.x = 1;
+      m_nctaid.y = 1;
+      m_nctaid.z = 1;
       m_gridid = 0;
       m_valid = true;
    }
-   void set_tid( int x, int y, int z)
-   {
-      m_tid[0] = x;
-      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;
-      m_ctaid[1] = y;
-      m_ctaid[2] = z;
-   }
-   void set_ntid( int x, int y, int z)
-   {
-      m_ntid[0] = x;
-      m_ntid[1] = y;
-      m_ntid[2] = z;
-   }
-   void set_nctaid( int x, int y, int z)
-   {
-      m_nctaid[0] = x;
-      m_nctaid[1] = y;
-      m_nctaid[2] = z;
-   }
+   void set_tid( dim3 tid ) { m_tid = tid; }
+   void cpy_tid_to_reg( dim3 tid );
+   void set_ctaid( dim3 ctaid ) { m_ctaid = ctaid; }
+   void set_ntid( dim3 tid ) { m_ntid = tid; }
+   void set_nctaid( dim3 cta_size ) { m_nctaid = cta_size; }
 
    unsigned get_builtin( int builtin_id, unsigned dim_mod ); 
 
@@ -384,7 +369,6 @@ public:
 
    unsigned next_instr()
    {
-      m_NPC = m_PC+1;   // increment to next instruction in case of no branch
       m_icount++;
       m_branch_taken = false;
       return m_PC;
@@ -404,6 +388,8 @@ public:
    void set_npc( const function_info *f );
    void callstack_push( unsigned npc, unsigned rpc, const symbol *return_var_src, const symbol *return_var_dst, unsigned call_uid );
    bool callstack_pop();
+   void callstack_push_plus( unsigned npc, unsigned rpc, const symbol *return_var_src, const symbol *return_var_dst, unsigned call_uid );
+   bool callstack_pop_plus();
    void dump_callstack() const;
    std::string get_location() const;
    const ptx_instruction *get_inst() const;
@@ -421,7 +407,7 @@ public:
    {
        return m_callstack.back().m_PC;
    }
-   void update_pc()
+   void update_pc( unsigned nbytes )
    {
       m_PC = m_NPC;
    }
@@ -451,10 +437,10 @@ private:
    unsigned m_uid;
    core_t *m_core;
    bool   m_valid;
-   unsigned m_ntid[3];
-   unsigned m_tid[3];
-   unsigned m_nctaid[3];
-   unsigned m_ctaid[3];
+   dim3   m_ntid;
+   dim3   m_tid;
+   dim3   m_nctaid;
+   dim3   m_ctaid;
    unsigned m_gridid;
    bool m_thread_done;
    unsigned m_hw_sid;