finished refactor+implement call/return passing (now working on simple example included in this CL)

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

5 files changed, 198 insertions, 149 deletions

diff --git a/benchmarks/CUDA/template/template_kernel.cu b/benchmarks/CUDA/template/template_kernel.cu
index 14610d8..adac694 100644
--- a/benchmarks/CUDA/template/template_kernel.cu
+++ b/benchmarks/CUDA/template/template_kernel.cu
@@ -33,7 +33,7 @@
 
 __device__ float foo( float bar )
 {
-	return 2.0f*bar;
+	return bar+1.0f;
 }
 
 __global__ void
@@ -51,7 +51,7 @@ testKernel( float* g_idata, float* g_odata)
   // read in input data from global memory
   // use the bank checker macro to check for bank conflicts during host
   // emulation
-  SDATA(tid) = foo(g_idata[tid])/2.0;
+  SDATA(tid) = foo(g_idata[tid])-1.0f;
   __syncthreads();
   printf("thread tid=%u reads %f from g_idata[]\n", tid, sdata[tid]);
 
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 0ec1896..fe3e1a1 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -687,141 +687,6 @@ extern int gpgpu_simd_model;
 void get_pdom_stack_top_info( unsigned sid, unsigned tid, unsigned *npc, unsigned *rpc );
 void gpgpusim_cuda_vprintf(const ptx_instruction * pI, const ptx_thread_info * thread, const function_info * target_func ); 
 
-
-class arg_buffer_t {
-public:
-   arg_buffer_t( const symbol *dst_sym, const operand_info &src_op, ptx_reg_t source_value ) : m_src_op(src_op)
-   {
-      m_dst = dst_sym;
-      if( dst_sym->is_reg() ) {
-         m_is_reg = true;
-         m_is_param = false;
-         assert( src_op.is_reg() );
-         m_reg_value = source_value;
-      } else {
-         m_is_param = true;
-         m_is_reg = false;
-         m_param_value = new ptx_reg_t(source_value);
-         m_param_bytes = sizeof(ptx_reg_t);
-      }
-   }
-   arg_buffer_t( const symbol *dst_sym, const operand_info &src_op, void *source_param_value_array, unsigned array_size ) : m_src_op(src_op)
-   {
-      m_dst = dst_sym;
-      if( dst_sym->is_reg() ) {
-         m_is_reg = true;
-         m_is_param = false;
-         assert( src_op.is_param_local() );
-         assert( dst_sym->get_size_in_bytes() == array_size );
-         switch( array_size ) {
-         case 1: m_reg_value.u8 = *(unsigned char*)source_param_value_array; break;
-         case 2: m_reg_value.u16 = *(unsigned short*)source_param_value_array; break;
-         case 4: m_reg_value.u32 = *(unsigned int*)source_param_value_array; break;
-         case 8: m_reg_value.u64 = *(unsigned long long*)source_param_value_array; break;
-         default:
-            printf("GPGPU-Sim PTX: ERROR ** source param size does not match known register sizes\n");
-            break;
-         }
-      } else {
-         // param
-         m_is_param = true;
-         m_is_reg = false;
-         m_param_value = new char[array_size];
-         m_param_bytes = array_size;
-         memcpy(m_param_value,source_param_value_array,array_size);
-      }
-   }
-
-   bool is_reg() const { return m_is_reg; }
-   ptx_reg_t get_reg() const 
-   { 
-      assert(m_is_reg); 
-      return m_reg_value; 
-   }
-
-   void *get_param_buffer() 
-   {
-      assert(m_is_param);
-      return m_param_value;
-   }
-   size_t get_param_buffer_size() const
-   {
-      assert(m_is_param);
-      return m_param_bytes;
-   }
-
-   const symbol *get_dst() const { return m_dst; }
-
-private:
-   // destination of copy
-   const symbol *m_dst;
-
-   // source operand
-   operand_info m_src_op;
-
-   // source information
-   bool m_is_reg;
-   bool m_is_param;
-
-   // source is register
-   ptx_reg_t m_reg_value;
-
-   // source is param
-   void     *m_param_value;
-   unsigned  m_param_bytes;
-};
-
-typedef std::list< arg_buffer_t > arg_buffer_list_t;
-
-void copy_arg_to_buffer(ptx_thread_info * thread, arg_buffer_list_t &arg_values, const operand_info &actual_param_op, const symbol * formal_param)
-{
-   if( actual_param_op.is_reg() )  {
-      ptx_reg_t value = thread->get_operand_value(actual_param_op);
-      arg_values.push_back( arg_buffer_t(formal_param,actual_param_op,value) );
-   } else {
-      assert( actual_param_op.is_param_local() ); 
-      unsigned size=formal_param->get_size_in_bytes();
-      addr_t from_addr = actual_param_op.get_symbol()->get_address();
-      char buffer[1024];
-      assert(size<1024); 
-      thread->m_local_mem->read(from_addr,size,buffer);
-      arg_values.push_back( arg_buffer_t(formal_param,actual_param_op,buffer,size) );
-   }
-}
-
-void copy_args_into_buffer_list( const ptx_instruction * pI, 
-                                 ptx_thread_info * thread, 
-                                 const function_info * target_func, 
-                                 arg_buffer_list_t &arg_values ) 
-{
-   unsigned n_return = target_func->has_return();
-   unsigned n_args = target_func->num_args();
-   for( unsigned arg=0; arg < n_args; arg ++ ) {
-      const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg);
-      const symbol *formal_param = target_func->get_arg(arg);
-      copy_arg_to_buffer(thread, arg_values, actual_param_op, formal_param);
-   }
-}
-
-void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &arg_values) 
-{
-   arg_buffer_list_t::iterator a;
-   for( a=arg_values.begin(); a != arg_values.end(); a++ ) {
-      if( a->is_reg() ) {
-         ptx_reg_t value = a->get_reg();
-         operand_info dst_reg = operand_info(a->get_dst()); 
-         thread->set_operand_value(dst_reg,value);
-      } else {  
-         void *buffer = a->get_param_buffer();
-         size_t size = a->get_param_buffer_size();
-         const symbol *dst = a->get_dst();
-         addr_t frame_offset = dst->get_address();
-         addr_t to_addr = thread->get_local_mem_stack_pointer() + frame_offset;
-         thread->m_local_mem->write(to_addr,size,buffer);
-      }
-   }
-}
-
 void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 
 {
    static unsigned call_uid_next = 1;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 1bf9c95..f51b022 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1360,3 +1360,61 @@ unsigned type_info_key::type_decode( int type, size_t &size, int &basic_type )
       return 0xDEADBEEF;
    }
 }
+
+arg_buffer_t copy_arg_to_buffer(ptx_thread_info * thread, operand_info actual_param_op, const symbol * formal_param)
+{
+   if( actual_param_op.is_reg() )  {
+      ptx_reg_t value = thread->get_operand_value(actual_param_op);
+      return arg_buffer_t(formal_param,actual_param_op,value);
+   } else if ( actual_param_op.is_param_local() ) {
+      unsigned size=formal_param->get_size_in_bytes();
+      addr_t frame_offset = actual_param_op.get_symbol()->get_address();
+      addr_t from_addr = thread->get_local_mem_stack_pointer() + frame_offset;
+      char buffer[1024];
+      assert(size<1024); 
+      thread->m_local_mem->read(from_addr,size,buffer);
+      return arg_buffer_t(formal_param,actual_param_op,buffer,size);
+   } else {
+      printf("GPGPU-Sim PTX: ERROR ** need to add support for this operand type in call/return\n");
+      abort();
+   }
+}
+
+void copy_args_into_buffer_list( const ptx_instruction * pI, 
+                                 ptx_thread_info * thread, 
+                                 const function_info * target_func, 
+                                 arg_buffer_list_t &arg_values ) 
+{
+   unsigned n_return = target_func->has_return();
+   unsigned n_args = target_func->num_args();
+   for( unsigned arg=0; arg < n_args; arg ++ ) {
+      const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg);
+      const symbol *formal_param = target_func->get_arg(arg);
+      arg_values.push_back( copy_arg_to_buffer(thread, actual_param_op, formal_param) );
+   }
+}
+
+void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a) 
+{
+   if( a.is_reg() ) {
+      ptx_reg_t value = a.get_reg();
+      operand_info dst_reg = operand_info(a.get_dst()); 
+      thread->set_operand_value(dst_reg,value);
+   } else {  
+      const void *buffer = a.get_param_buffer();
+      size_t size = a.get_param_buffer_size();
+      const symbol *dst = a.get_dst();
+      addr_t frame_offset = dst->get_address();
+      addr_t to_addr = thread->get_local_mem_stack_pointer() + frame_offset;
+      thread->m_local_mem->write(to_addr,size,buffer);
+   }
+}
+
+void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &arg_values) 
+{
+   arg_buffer_list_t::iterator a;
+   for( a=arg_values.begin(); a != arg_values.end(); a++ ) {
+      copy_buffer_to_frame(thread, *a);
+   }
+}
+
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 4361cb0..0c63e41 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1070,6 +1070,132 @@ private:
    static unsigned sm_next_uid;
 };
 
+class arg_buffer_t {
+public:
+   arg_buffer_t()
+   {
+      m_is_reg=false;
+      m_is_param=false;
+      m_param_value=NULL;
+   }
+   arg_buffer_t( const arg_buffer_t &another )
+   {
+      make_copy(another);
+   }
+   void make_copy( const arg_buffer_t &another )
+   {
+      m_dst = another.m_dst;
+      m_src_op = another.m_src_op;
+      m_is_reg = another.m_is_reg;
+      m_is_param = another.m_is_param;
+      m_reg_value = another.m_reg_value;
+      m_param_bytes = another.m_param_bytes;
+      if( m_is_param ) {
+         m_param_value = malloc(m_param_bytes);
+         memcpy(m_param_value,another.m_param_value,m_param_bytes);
+      }
+   }
+   void operator=( const arg_buffer_t &another )
+   {
+      make_copy(another);
+   }
+   ~arg_buffer_t()
+   {
+      if( m_is_param ) 
+         free(m_param_value);
+   }
+   arg_buffer_t( const symbol *dst_sym, const operand_info &src_op, ptx_reg_t source_value ) : m_src_op(src_op)
+   {
+      m_dst = dst_sym;
+      if( dst_sym->is_reg() ) {
+         m_is_reg = true;
+         m_is_param = false;
+         assert( src_op.is_reg() );
+         m_reg_value = source_value;
+      } else {
+         m_is_param = true;
+         m_is_reg = false;
+         m_param_value = calloc(sizeof(ptx_reg_t),1);
+         //new (m_param_value) ptx_reg_t(source_value);
+         memcpy(m_param_value,&source_value,sizeof(ptx_reg_t));
+         m_param_bytes = sizeof(ptx_reg_t);
+      }
+   }
+   arg_buffer_t( const symbol *dst_sym, const operand_info &src_op, void *source_param_value_array, unsigned array_size ) : m_src_op(src_op)
+   {
+      m_dst = dst_sym;
+      if( dst_sym->is_reg() ) {
+         m_is_reg = true;
+         m_is_param = false;
+         assert( src_op.is_param_local() );
+         assert( dst_sym->get_size_in_bytes() == array_size );
+         switch( array_size ) {
+         case 1: m_reg_value.u8 = *(unsigned char*)source_param_value_array; break;
+         case 2: m_reg_value.u16 = *(unsigned short*)source_param_value_array; break;
+         case 4: m_reg_value.u32 = *(unsigned int*)source_param_value_array; break;
+         case 8: m_reg_value.u64 = *(unsigned long long*)source_param_value_array; break;
+         default:
+            printf("GPGPU-Sim PTX: ERROR ** source param size does not match known register sizes\n");
+            break;
+         }
+      } else {
+         // param
+         m_is_param = true;
+         m_is_reg = false;
+         m_param_value = calloc(array_size,1);
+         m_param_bytes = array_size;
+         memcpy(m_param_value,source_param_value_array,array_size);
+      }
+   }
+
+   bool is_reg() const { return m_is_reg; }
+   ptx_reg_t get_reg() const 
+   { 
+      assert(m_is_reg); 
+      return m_reg_value; 
+   }
+
+   const void *get_param_buffer() const
+   {
+      assert(m_is_param);
+      return m_param_value;
+   }
+   size_t get_param_buffer_size() const
+   {
+      assert(m_is_param);
+      return m_param_bytes;
+   }
+
+   const symbol *get_dst() const { return m_dst; }
+
+private:
+   // destination of copy
+   const symbol *m_dst;
+
+   // source operand
+   operand_info m_src_op;
+
+   // source information
+   bool m_is_reg;
+   bool m_is_param;
+
+   // source is register
+   ptx_reg_t m_reg_value;
+
+   // source is param
+   void     *m_param_value;
+   unsigned  m_param_bytes;
+};
+
+typedef std::list< arg_buffer_t > arg_buffer_list_t;
+arg_buffer_t copy_arg_to_buffer(ptx_thread_info * thread, operand_info actual_param_op, const symbol * formal_param);
+void copy_args_into_buffer_list( const ptx_instruction * pI, 
+                                 ptx_thread_info * thread, 
+                                 const function_info * target_func, 
+                                 arg_buffer_list_t &arg_values ); 
+void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &arg_values);
+void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a);
+
 /*******************************/
 // These declarations should be identical to those in ./../../cuda-sim-dev/libcuda/texture_types.h
 enum cudaChannelFormatKind {
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index cf0236a..3974408 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -369,29 +369,29 @@ extern void set_operand_value( const symbol *dst, const ptx_reg_t &data );
 bool ptx_thread_info::callstack_pop()
 {
    assert( !m_callstack.empty() );
-   assert( m_local_mem_stack_pointer >= m_func_info->local_mem_framesize() );
-   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;
+   assert( m_local_mem_stack_pointer >= m_callstack.back().m_func_info->local_mem_framesize() );
    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)) ); 
+   assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value
 
-   // TODO: read param local memory source
-   ptx_reg_t rval;
+   // read return value from callee frame
+   arg_buffer_t buffer;
    if( rv_src != NULL ) 
-      rval = get_operand_value(rv_src);
+      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;
    m_local_mem_stack_pointer -= m_func_info->local_mem_framesize(); 
    m_callstack.pop_back();
    m_regs.pop_back();
 
-   // TODO: write param local memory destination
+   // write return value into caller frame
    if( rv_dst != NULL ) 
-      set_operand_value(rv_dst,rval);
+      copy_buffer_to_frame(this, buffer);
 
    return m_callstack.empty();
 }