- parameter memory and active threads now part of kernel_info_t:

  Parameters are finalized at kernel launch, which means the contents
  of parameter memory are initialized. Kernel arguement names have a
  fixed order, hence same address should be assigned on subsequent
  kernel launches of same kernel in other streams provided the data size
  param_t::size of arguments for each kernel launch is identical (an
  assertion has been added to check this is true).

- passing regression

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

16 files changed, 125 insertions, 100 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index f5e7810..4d5d856 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -139,7 +139,7 @@ extern void exit_simulation();
 static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu );
 static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu );
 
-static kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, 
+static 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,
@@ -909,8 +909,8 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
    cudaStream_t stream = config.get_stream();
    printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
           g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
-   kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
-   std::string kname = grid.name();
+   kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
+   std::string kname = grid->name();
    dim3 gridDim = config.grid_dim();
    dim3 blockDim = config.block_dim();
    printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
@@ -1559,13 +1559,14 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu
    return nc_bytes;
 }
 
-kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, 
+kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, 
                                             gpgpu_ptx_sim_arg_list_t args, 
                                             struct dim3 gridDim, 
                                             struct dim3 blockDim, 
                                             CUctx_st* context )
 {
    function_info *entry = context->get_kernel(hostFun);
+   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
    if( entry == NULL ) {
        printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found\n");
        abort();
@@ -1577,10 +1578,10 @@ kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
       argn++;
    }
 
-   entry->finalize(context->get_device()->get_gpgpu()->get_param_memory());
+   entry->finalize(result->get_param_memory());
    g_ptx_kernel_count++; 
    fflush(stdout);
 
-   return kernel_info_t(gridDim,blockDim,entry);
+   return result;
 }
 
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 8943061..7a23f37 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -873,7 +873,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
    gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
 
-   kernel_info_t grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
+   kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
    if ( g_ptx_sim_mode )
       gpgpu_opencl_ptx_sim_main_func( grid );
    else
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index f779e3d..e883455 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -52,7 +52,6 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config )
     : m_function_model_config(config)
 {
    m_global_mem = new memory_space_impl<8192>("global",64*1024);
-   m_param_mem = new memory_space_impl<8192>("param",64*1024);
    m_tex_mem = new memory_space_impl<8192>("tex",64*1024);
    m_surf_mem = new memory_space_impl<8192>("surf",64*1024);
 
@@ -310,6 +309,25 @@ void warp_inst_t::generate_mem_accesses()
 
 unsigned kernel_info_t::m_next_uid = 1;
 
+kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry )
+{
+    m_kernel_entry=entry;
+    m_grid_dim=gridDim;
+    m_block_dim=blockDim;
+    m_next_cta.x=0;
+    m_next_cta.y=0;
+    m_next_cta.z=0;
+    m_next_tid=m_next_cta;
+    m_num_cores_running=0;
+    m_uid = m_next_uid++;
+    m_param_mem = new memory_space_impl<8192>("param",64*1024);
+}
+
+kernel_info_t::~kernel_info_t()
+{
+    assert( m_active_threads.empty() );
+    delete m_param_mem;
+}
 
 std::string kernel_info_t::name() const
 {
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 77a35a7..2332670 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -87,26 +87,16 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound);
 
 class kernel_info_t {
 public:
-   kernel_info_t()
-   {
-      m_valid=false;
-      m_kernel_entry=NULL;
-      m_uid=0;
-      m_num_cores_running=0;
-   }
-   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry )
-   {
-      m_valid=true;
-      m_kernel_entry=entry;
-      m_grid_dim=gridDim;
-      m_block_dim=blockDim;
-      m_next_cta.x=0;
-      m_next_cta.y=0;
-      m_next_cta.z=0;
-      m_next_tid=m_next_cta;
-      m_num_cores_running=0;
-      m_uid = m_next_uid++;
-   }
+//   kernel_info_t()
+//   {
+//      m_valid=false;
+//      m_kernel_entry=NULL;
+//      m_uid=0;
+//      m_num_cores_running=0;
+//      m_param_mem=NULL;
+//   }
+   kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry );
+   ~kernel_info_t();
 
    void inc_running() { m_num_cores_running++; }
    void dec_running()
@@ -115,10 +105,9 @@ public:
        m_num_cores_running--; 
    }
    bool running() const { return m_num_cores_running>0; }
-   bool valid() const  { return m_valid; }
    bool done() const 
    {
-       return (!valid()) || (no_more_ctas_to_run() && !running());
+       return no_more_ctas_to_run() && !running();
    }
    class function_info *entry() { return m_kernel_entry; }
    const class function_info *entry() const { return m_kernel_entry; }
@@ -162,8 +151,13 @@ public:
    unsigned get_uid() const { return m_uid; }
    std::string name() const;
 
+   std::list<class ptx_thread_info *> &active_threads() { return m_active_threads; }
+   class memory_space *get_param_memory() { return m_param_mem; }
+
 private:
-   bool m_valid;
+   kernel_info_t( const kernel_info_t & ); // disable copy constructor
+   void operator=( const kernel_info_t & ); // disable copy operator
+
    class function_info *m_kernel_entry;
 
    unsigned m_uid;
@@ -175,6 +169,9 @@ private:
    dim3 m_next_tid;
 
    unsigned m_num_cores_running;
+
+   std::list<class ptx_thread_info *> m_active_threads;
+   class memory_space *m_param_mem;
 };
 
 struct core_config {
@@ -319,7 +316,6 @@ public:
     class memory_space *get_global_memory() { return m_global_mem; }
     class memory_space *get_tex_memory() { return m_tex_mem; }
     class memory_space *get_surf_memory() { return m_surf_mem; }
-    class memory_space *get_param_memory() { return m_param_mem; }
 
     void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* texref, const struct cudaArray* array);
     void gpgpu_ptx_sim_bindNameToTexture(const char* name, const struct textureReference* texref);
@@ -354,7 +350,6 @@ protected:
     class memory_space *m_global_mem;
     class memory_space *m_tex_mem;
     class memory_space *m_surf_mem;
-    class memory_space *m_param_mem;
     
     unsigned long long m_dev_malloc;
     
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index b6a15e8..e33181f 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1075,7 +1075,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
                               unsigned hw_warp_id,
                               gpgpu_t *gpu )
 {
-   static std::list<ptx_thread_info *> active_threads;
+   std::list<ptx_thread_info *> &active_threads = kernel.active_threads();
+
    static std::map<unsigned,memory_space*> shared_memory_lookup;
    static std::map<unsigned,ptx_cta_info*> ptx_cta_lookup;
    static std::map<unsigned,std::map<unsigned,memory_space*> > local_memory_lookup;
@@ -1096,7 +1097,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
       *thread_info = NULL;
    }
 
-   if ( !active_threads.empty() ) { //if g_active_threads not empty...
+   if ( !active_threads.empty() ) {
       assert( active_threads.size() <= threads_left );
       ptx_thread_info *thd = active_threads.front(); 
       active_threads.pop_front();
@@ -1156,7 +1157,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
       dim3 tid3d = kernel.get_next_thread_id_3d();
       kernel.increment_thread_id();
       new_tid += tid;
-      ptx_thread_info *thd = new ptx_thread_info();
+      ptx_thread_info *thd = new ptx_thread_info(kernel);
    
       memory_space *local_mem = NULL;
       std::map<unsigned,memory_space*>::iterator l = local_mem_lookup.find(new_tid);
@@ -1210,23 +1211,24 @@ size_t get_kernel_code_size( class function_info *entry )
 }
 
 
-kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+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,
                                              gpgpu_t *gpu )
 {
+   kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
    unsigned argcount=args.size();
    unsigned argn=1;
    for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
       entry->add_param_data(argcount-argn,&(*a));
       argn++;
    }
-   entry->finalize(gpu->get_param_memory());
+   entry->finalize(result->get_param_memory());
    g_ptx_kernel_count++; 
    fflush(stdout);
 
-   return kernel_info_t(gridDim,blockDim,entry);
+   return result;
 }
 
 const char *g_gpgpusim_version_string = "2.1.1b (beta)";
@@ -1395,7 +1397,7 @@ ptx_cta_info *g_func_cta_info = NULL;
 
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
-void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel )
+void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel )
 {
    printf("GPGPU-Sim: Performing Functional Simulation...\n");
 
@@ -1416,7 +1418,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel )
             g_func_cta_info->check_cta_thread_status_and_reset();
             while( kernel.more_threads_in_cta() ) {
                      memory_space *local_mem = NULL;
-                     ptx_thread_info *thd = new ptx_thread_info();
+                     ptx_thread_info *thd = new ptx_thread_info(kernel);
                      dim3 tid3d = kernel.get_next_thread_id_3d();
                      unsigned lm_idx = kernel.get_next_thread_id();
                      kernel.increment_thread_id();
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index ef045fb..3bc915d 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -20,12 +20,12 @@ extern void ** g_inst_op_classification_stat;
 extern int g_ptx_kernel_count; // used for classification stat collection purposes 
 
 
-extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+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, 
                                                           class gpgpu_t *gpu );
-extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel );
+extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel );
 extern void   print_splash();
 extern void   gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size );
 extern void   gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size );
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 59f4522..3ab0594 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -211,7 +211,8 @@ ptx_thread_info::~ptx_thread_info()
    g_ptx_thread_info_delete_count++;
 }
 
-ptx_thread_info::ptx_thread_info()
+ptx_thread_info::ptx_thread_info( kernel_info_t &kernel )
+    : m_kernel(kernel)
 {
    m_uid = g_ptx_thread_info_uid_next++;
    m_core = NULL;
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index d60955b..cc555cf 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -274,7 +274,7 @@ private:
 class ptx_thread_info {
 public:
    ~ptx_thread_info();
-   ptx_thread_info();
+   ptx_thread_info( kernel_info_t &kernel );
 
    void init(gpgpu_t *gpu, core_t *core, unsigned sid, unsigned cta_id, unsigned wid, unsigned tid ) 
    { 
@@ -429,7 +429,7 @@ public:
    memory_space *get_global_memory() { return m_gpu->get_global_memory(); }
    memory_space *get_tex_memory() { return m_gpu->get_tex_memory(); }
    memory_space *get_surf_memory() { return m_gpu->get_surf_memory(); }
-   memory_space *get_param_memory() { return m_gpu->get_param_memory(); }
+   memory_space *get_param_memory() { return m_kernel.get_param_memory(); }
    const gpgpu_functional_sim_config &get_config() const { return m_gpu->get_config(); }
 
 public:
@@ -446,6 +446,7 @@ private:
 
 
    unsigned m_uid;
+   kernel_info_t &m_kernel;
    core_t *m_core;
    gpgpu_t *m_gpu;
    bool   m_valid;
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 9ae05d8..05a9b55 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -309,9 +309,9 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound)
    }
 }
 
-void gpgpu_sim::launch( kernel_info_t &kinfo )
+void gpgpu_sim::launch( kernel_info_t *kinfo )
 {
-   unsigned cta_size = kinfo.threads_per_cta();
+   unsigned cta_size = kinfo->threads_per_cta();
    if ( cta_size > m_shader_config->n_thread_per_shader ) {
       printf("Execution error: Shader kernel CTA (block) size is too large for microarch config.\n");
       printf("                 CTA size (x*y*z) = %u, max supported = %u\n", cta_size, 
@@ -322,7 +322,7 @@ void gpgpu_sim::launch( kernel_info_t &kinfo )
    }
    unsigned n=0;
    for(n=0; n < m_running_kernels.size(); n++ ) {
-       if( m_running_kernels[n].done() ) {
+       if( (NULL==m_running_kernels[n]) || m_running_kernels[n]->done() ) {
            m_running_kernels[n] = kinfo;
            break;
        }
@@ -333,7 +333,7 @@ void gpgpu_sim::launch( kernel_info_t &kinfo )
 bool gpgpu_sim::can_start_kernel()
 {
    for(unsigned n=0; n < m_running_kernels.size(); n++ ) {
-       if( m_running_kernels[n].done() ) 
+       if( (NULL==m_running_kernels[n]) || m_running_kernels[n]->done() ) 
            return true;
    }
    return false;
@@ -346,7 +346,7 @@ bool gpgpu_sim::get_more_cta_left() const
           return false;
    }
    for(unsigned n=0; n < m_running_kernels.size(); n++ ) {
-       if( m_running_kernels[n].valid() && !m_running_kernels[n].no_more_ctas_to_run() ) 
+       if( m_running_kernels[n] && !m_running_kernels[n]->no_more_ctas_to_run() ) 
            return true;
    }
    return false;
@@ -356,9 +356,9 @@ kernel_info_t *gpgpu_sim::select_kernel()
 {
     for(unsigned n=0; n < m_running_kernels.size(); n++ ) {
         unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel;
-        if( m_running_kernels[idx].valid() && !m_running_kernels[idx].no_more_ctas_to_run() ) {
+        if( m_running_kernels[idx] && !m_running_kernels[idx]->no_more_ctas_to_run() ) {
             m_last_issued_kernel=idx;
-            return &m_running_kernels[idx];
+            return m_running_kernels[idx];
         }
     }
     return NULL;
@@ -373,6 +373,20 @@ unsigned gpgpu_sim::finished_kernel()
     return result;
 }
 
+void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) 
+{ 
+    unsigned uid = kernel->get_uid();
+    m_finished_kernel.push_back(uid);
+    std::vector<kernel_info_t*>::iterator k;
+    for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) {
+        if( *k == kernel ) {
+            *k = NULL;
+            break;
+        }
+    }
+    assert( k != m_running_kernels.end() ); 
+}
+
 void set_ptx_warp_size(const struct core_config * warp_size);
 
 gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) 
@@ -404,7 +418,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config )
     time_vector_create(NUM_MEM_REQ_STAT);
     fprintf(stdout, "GPGPU-Sim uArch: performance model initialization complete.\n");
 
-    m_running_kernels.resize( config.max_concurrent_kernel );
+    m_running_kernels.resize( config.max_concurrent_kernel, NULL );
     m_last_issued_kernel = 0;
     m_last_cluster_issue = 0;
 }
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 5631ab9..f229965 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -242,10 +242,10 @@ public:
 
    void set_prop( struct cudaDeviceProp *prop );
 
-   void launch( kernel_info_t &kinfo );
+   void launch( kernel_info_t *kinfo );
    bool can_start_kernel();
    unsigned finished_kernel();
-   void set_kernel_done( unsigned uid ) { m_finished_kernel.push_back(uid); }
+   void set_kernel_done( kernel_info_t *kernel );
 
    void init();
    void cycle();
@@ -289,7 +289,7 @@ private:
    class simt_core_cluster **m_cluster;
    class memory_partition_unit **m_memory_partition_unit;
 
-   std::vector<kernel_info_t> m_running_kernels;
+   std::vector<kernel_info_t*> m_running_kernels;
    unsigned m_last_issued_kernel;
 
    std::list<unsigned> m_finished_kernel;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 07aeb6d..2380bcf 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1207,11 +1207,12 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num )
       if( m_n_active_cta == 0 ) {
           assert( m_kernel != NULL );
           m_kernel->dec_running();
-          printf("GPGPU-Sim uArch: Shader %u empty.\n", m_sid );
+          printf("GPGPU-Sim uArch: Shader %u empty (release kernel %u \'%s\').\n", m_sid, m_kernel->get_uid(),
+                 m_kernel->name().c_str() );
           if( m_kernel->no_more_ctas_to_run() ) {
               if( !m_kernel->running() ) {
-                  m_gpu->set_kernel_done( m_kernel->get_uid() );
                   printf("GPGPU-Sim uArch: GPU detected kernel \'%s\' finished on shader %u.\n", m_kernel->name().c_str(), m_sid );
+                  m_gpu->set_kernel_done( m_kernel );
               }
           }
           m_kernel=NULL;
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 7f41cfe..2b834a6 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1123,8 +1123,9 @@ public:
     {
         assert(k);
         m_kernel=k; 
-        if(k) k->inc_running(); 
-        printf("GPGPU-Sim uArch: Shader %d kernel = 0x%p\n", m_sid, m_kernel );
+        k->inc_running(); 
+        printf("GPGPU-Sim uArch: Shader %d bind to kernel %u \'%s\'\n", m_sid, m_kernel->get_uid(),
+                 m_kernel->name().c_str() );
     }
     
     // accessors
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 00ba5a2..cf70f24 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -137,6 +137,7 @@ void *gpgpu_sim_thread_concurrent(void*)
         g_sim_active = true;
         pthread_mutex_unlock(&g_sim_lock);
         bool active = false;
+        bool sim_cycles = false;
         do {
             // check if a kernel has completed
             unsigned grid_uid = g_the_gpu->finished_kernel();
@@ -148,14 +149,17 @@ void *gpgpu_sim_thread_concurrent(void*)
             op.do_operation(g_the_gpu);
     
             // simulate a clock cycle on the GPU 
-            if( g_the_gpu->active() ) 
+            if( g_the_gpu->active() ) { 
                 g_the_gpu->cycle();
+                sim_cycles = true;
+            }
             g_the_gpu->deadlock_check();
             active = g_the_gpu->active() || !g_stream_manager->empty();
         } while( active );
         printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n");
         fflush(stdout);
-        g_the_gpu->print_stats();
+        if( sim_cycles ) 
+            g_the_gpu->print_stats();
         pthread_mutex_lock(&g_sim_lock);
         g_sim_active = false;
         pthread_mutex_unlock(&g_sim_lock);
@@ -254,15 +258,7 @@ void print_simulation_time()
    fflush(stdout);
 }
 
-int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid )
-{
-   g_the_gpu->launch(grid);
-   //sem_post(&g_sim_signal_start);
-   //sem_wait(&g_sim_signal_finish);
-   return 0;
-}
-
-int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid )
+int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid )
 {
    g_the_gpu->launch(grid);
    sem_post(&g_sim_signal_start);
@@ -270,7 +266,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid )
    return 0;
 }
 
-int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid )
+int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid )
 {
    printf("GPGPU-Sim PTX API: OpenCL functional-only simulation not yet implemented (use performance simulation)\n");
    exit(1);
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index b310752..4f1d503 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -73,8 +73,7 @@ extern time_t g_simulation_starttime;
 class gpgpu_sim *gpgpu_ptx_sim_init_perf();
 void start_sim_thread(int api);
 
-int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid );
-int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid );
-int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid );
+int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid );
+int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid );
 
 #endif
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 5d1cfe2..54e2ee7 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -182,9 +182,9 @@ void stream_operation::do_operation( gpgpu_sim *gpu )
         break;
     case stream_kernel_launch:
         if( gpu->can_start_kernel() ) {
-            printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel.name().c_str() );
+            printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() );
             if( m_sim_mode )
-                gpgpu_cuda_ptx_sim_main_func( m_kernel );
+                gpgpu_cuda_ptx_sim_main_func( *m_kernel );
             else
                 gpu->launch( m_kernel );
         }
@@ -231,9 +231,11 @@ void stream_manager::register_finished_kernel( unsigned grid_uid )
     // called by gpu simulation thread
     pthread_mutex_lock(&m_lock);
     CUstream_st *stream = m_grid_id_to_stream[grid_uid];
-    assert( grid_uid == stream->front().get_kernel().get_uid() );
+    kernel_info_t *kernel = stream->front().get_kernel();
+    assert( grid_uid == kernel->get_uid() );
     stream->record_next_done();
     m_grid_id_to_stream.erase(grid_uid);
+    delete kernel;
     pthread_mutex_unlock(&m_lock);
 }
 
@@ -249,7 +251,7 @@ stream_operation stream_manager::front()
             if( !m_stream_zero.busy() ) {
                 result = m_stream_zero.next();
                 if( result.is_kernel() ) {
-                    unsigned grid_id = result.get_kernel().get_uid();
+                    unsigned grid_id = result.get_kernel()->get_uid();
                     m_grid_id_to_stream[grid_id] = &m_stream_zero;
                 }
             }
@@ -262,18 +264,10 @@ stream_operation stream_manager::front()
             CUstream_st *stream = *s;
             if( !stream->busy() && !stream->empty() ) {
                 result = stream->next();
-//                stream_barrier *b = result.get_barrier();
-//                if( b && b->value() > 1 ) {
-//                    b->dec();
-//                    result = stream_operation();
-//                } else {
-//                    assert( b->value() == 1 );
-//                    delete b;
-                    if( result.is_kernel() ) {
-                        unsigned grid_id = result.get_kernel().get_uid();
-                        m_grid_id_to_stream[grid_id] = stream;
-                    }
-//                }
+                if( result.is_kernel() ) {
+                    unsigned grid_id = result.get_kernel()->get_uid();
+                    m_grid_id_to_stream[grid_id] = stream;
+                }
                 break;
             }
         }
diff --git a/src/stream_manager.h b/src/stream_manager.h
index d47f0c7..ed3bc9b 100644
--- a/src/stream_manager.h
+++ b/src/stream_manager.h
@@ -95,23 +95,25 @@ class stream_operation {
 public:
     stream_operation()
     {
+        m_kernel=NULL;
         m_type = stream_no_op;
         m_stream = NULL;
         m_done=true;
     }
     stream_operation( const void *src, const char *symbol, size_t count, size_t offset, struct CUstream_st *stream )
     {
+        m_kernel=NULL;
         m_stream = stream;
         m_type=stream_memcpy_to_symbol;
         m_host_address_src=src;
         m_symbol=symbol;
         m_cnt=count;
         m_offset=offset;
-
         m_done=false;
     }
     stream_operation( const char *symbol, void *dst, size_t count, size_t offset, struct CUstream_st *stream )
     {
+        m_kernel=NULL;
         m_stream = stream;
         m_type=stream_memcpy_from_symbol;
         m_host_address_dst=dst;
@@ -120,7 +122,7 @@ public:
         m_offset=offset;
         m_done=false;
     }
-    stream_operation( const kernel_info_t &kernel, bool sim_mode, struct CUstream_st *stream )
+    stream_operation( kernel_info_t *kernel, bool sim_mode, struct CUstream_st *stream )
     {
         m_type=stream_kernel_launch;
         m_kernel=kernel;
@@ -130,6 +132,7 @@ public:
     }
     stream_operation( class CUevent_st *e, struct CUstream_st *stream )
     {
+        m_kernel=NULL;
         m_type=stream_event;
         m_event=e;
         m_stream=stream;
@@ -137,6 +140,7 @@ public:
     }
     stream_operation( const void *host_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream )
     {
+        m_kernel=NULL;
         m_type=stream_memcpy_host_to_device;
         m_host_address_src  =host_address_src;
         m_device_address_dst=device_address_dst;
@@ -149,6 +153,7 @@ public:
     }
     stream_operation( size_t device_address_src, void *host_address_dst, size_t cnt, struct CUstream_st *stream  )
     {
+        m_kernel=NULL;
         m_type=stream_memcpy_device_to_host;
         m_device_address_src=device_address_src;
         m_host_address_dst=host_address_dst;
@@ -161,6 +166,7 @@ public:
     }
     stream_operation( size_t device_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream  )
     {
+        m_kernel=NULL;
         m_type=stream_memcpy_device_to_device;
         m_device_address_src=device_address_src;
         m_device_address_dst=device_address_dst;
@@ -172,8 +178,6 @@ public:
         m_done=false;
     }
 
-    //void add_barrier( stream_barrier *b ) { m_barrier=b; }
-    //stream_barrier *get_barrier() { return m_barrier; }
     bool is_kernel() const { return m_type == stream_kernel_launch; }
     bool is_mem() const {
         return m_type == stream_memcpy_host_to_device ||
@@ -182,7 +186,7 @@ public:
     }
     bool is_noop() const { return m_type == stream_no_op; }
     bool is_done() const { return m_done; }
-    kernel_info_t &get_kernel() { return m_kernel; }
+    kernel_info_t *get_kernel() { return m_kernel; }
     void do_operation( gpgpu_sim *gpu );
     void print( FILE *fp ) const;
     struct CUstream_st *get_stream() { return m_stream; }
@@ -204,10 +208,8 @@ private:
     size_t m_offset;
 
     bool m_sim_mode;
-    kernel_info_t m_kernel;
+    kernel_info_t *m_kernel;
     class CUevent_st *m_event;
-
-    //stream_barrier *m_barrier;
 };
 
 class CUevent_st {