- Checkpointing new support for concurrent kernel execution (CUDA only, not OpenCL)

  This changelist adds full support for streams supported by a new class,
  stream_manager and enables concurrent execution of kernels from different
  streams.

- fast_regression.sh fails for simpleMultiCopy, simpleStreams (other tests
  passing)

** Known issues **
- Kernel parameter passing is not done correctly for concurrent kernel execution
  (somehow concurrentKernels is not affected by this): the parameters are
  stored inside function_info, which is shared among parallel kernel launches
  so that the values passed into the launch are likely to get overwritten if
  multiple grids are launched in parallel streams.

- Statistics are printed out whenever the simulation thread runs out of
  cuda commands (doesn't make sense to print out when a kernel ends during
  concurrent kernel execution).  This will probably require further tweaking
  so as to be more compatible with data collection scripts.

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

16 files changed, 1205 insertions, 317 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 1990076..f5e7810 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -128,6 +128,13 @@
 #include "../src/cuda-sim/ptx_ir.h"
 #include "../src/cuda-sim/ptx_parser.h"
 #include "../src/gpgpusim_entrypoint.h"
+#include "../src/stream_manager.h"
+
+#include <pthread.h>
+#include <semaphore.h>
+
+extern void synchronize();
+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 );
@@ -161,6 +168,8 @@ struct cudaArray
 
 cudaError_t g_last_cudaError = cudaSuccess;
 
+extern stream_manager *g_stream_manager;
+
 void register_ptx_function( const char *name, function_info *impl )
 {
    // no longer need this 
@@ -256,6 +265,32 @@ private:
    std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
 };
 
+class kernel_config {
+public:
+    kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
+    { 
+        m_GridDim=GridDim; 
+        m_BlockDim=BlockDim; 
+        m_sharedMem=sharedMem;
+        m_stream = stream;
+    }
+    void set_arg( const void *arg, size_t size, size_t offset )
+    {
+        m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
+    }
+    dim3 grid_dim() const { return m_GridDim; }
+    dim3 block_dim() const { return m_BlockDim; }
+    gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
+    struct CUstream_st *get_stream() { return m_stream; }
+    
+private:
+    dim3 m_GridDim;
+    dim3 m_BlockDim;
+    size_t m_sharedMem;
+    struct CUstream_st *m_stream;
+    gpgpu_ptx_sim_arg_list_t m_args;
+};
+
 class _cuda_device_id *GPGPUSim_Init()
 {
    static _cuda_device_id *the_device = NULL;
@@ -287,6 +322,7 @@ class _cuda_device_id *GPGPUSim_Init()
       the_gpu->set_prop(prop);
       the_device = new _cuda_device_id(the_gpu);
    } 
+   start_sim_thread(1);
    return the_device;
 }
 
@@ -353,67 +389,6 @@ void gpgpusim_ptx_assert_impl( int test_value, const char *func, const char *fil
       gpgpusim_ptx_error_impl(func, file, line, msg);
 }
 
-class kernel_config {
-public:
-   kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, cudaStream_t stream ) 
-   { 
-      m_GridDim=GridDim; 
-      m_BlockDim=BlockDim; 
-      m_sharedMem=sharedMem;
-      m_stream=stream;
-   }
-   void set_arg(const void *arg, size_t size, size_t offset)
-   {
-      m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
-   }
-   dim3 grid_dim() const { return m_GridDim; }
-   dim3 block_dim() const { return m_BlockDim; }
-   gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
-    
-private:
-   dim3 m_GridDim;
-   dim3 m_BlockDim;
-   size_t m_sharedMem;
-   cudaStream_t m_stream;
-   gpgpu_ptx_sim_arg_list_t m_args;
-};
-
-struct CUstream_st {
-   int _;
-};
-
-class CUevent_st {
-public:
-   CUevent_st( bool blocking )
-   {
-      m_uid = ++m_next_event_uid;
-      m_blocking = blocking;
-      m_updates = 0;
-      m_wallclock = 0;
-      m_gpu_tot_sim_cycle = 0;
-      m_done = true;
-   }
-   void update( double cycle, time_t clk )
-   {
-      m_updates++;
-      m_wallclock=clk;
-      m_gpu_tot_sim_cycle=cycle;
-   }
-   //void set_done() { assert(!m_done); m_done=true; }
-   int get_uid() const { return m_uid; }
-   unsigned num_updates() const { return m_updates; }
-   bool done() const { return m_done; }
-   time_t clock() const { return m_wallclock; }
-private:
-   int m_uid;
-   bool m_blocking;
-   bool m_done;
-   int m_updates;
-   time_t m_wallclock;
-   double m_gpu_tot_sim_cycle;
-
-   static int m_next_event_uid;
-};
 
 typedef std::map<unsigned,CUevent_st*> event_tracker_t;
 
@@ -516,17 +491,17 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
 *                                                                              *
 *******************************************************************************/
 
- __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
+__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
 {
-   CUctx_st *context = GPGPUSim_Context();
-   gpgpu_t *gpu = context->get_device()->get_gpgpu();
+   //CUctx_st *context = GPGPUSim_Context();
+   //gpgpu_t *gpu = context->get_device()->get_gpgpu();
    printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
    if( kind == cudaMemcpyHostToDevice ) 
-      gpu->memcpy_to_gpu( (size_t)dst, src, count ); 
+       g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
    else if( kind == cudaMemcpyDeviceToHost ) 
-     	gpu->memcpy_from_gpu( dst, (size_t)src, count ); 
+       g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
    else if( kind == cudaMemcpyDeviceToDevice ) 
-      gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count ); 
+       g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
    else {
       printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); 
       abort();
@@ -636,20 +611,23 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
 
  __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
 {
-   CUctx_st *context = GPGPUSim_Context();
+   //CUctx_st *context = GPGPUSim_Context();
    assert(kind == cudaMemcpyHostToDevice);
    printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
-   gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
+    //stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
+   g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
+   //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
    return g_last_cudaError = cudaSuccess;
 }
 
 
  __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost))
 {
-   CUctx_st *context = GPGPUSim_Context();
+   //CUctx_st *context = GPGPUSim_Context();
    assert(kind == cudaMemcpyDeviceToHost);
    printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
-   gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
+   g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) );
+   //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
    return g_last_cudaError = cudaSuccess;
 }
 
@@ -663,8 +641,14 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
 
  __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
 {
-   printf("GPGPU-Sim PTX: warning cudaMemcpyAsync is implemented as blocking in this version of GPGPU-Sim...\n");
-   return cudaMemcpy(dst,src,count,kind);
+    switch( kind ) {
+    case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,stream) ); break;
+    case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,stream) ); break;
+    case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,stream) ); break;
+    default:
+        abort();
+    }
+    return g_last_cudaError = cudaSuccess;
 }
 
 
@@ -920,15 +904,20 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
    char *mode = getenv("PTX_SIM_MODE_FUNC");
    if( mode ) 
       sscanf(mode,"%u", &g_ptx_sim_mode);
-   printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s)\n", hostFun,
-          g_ptx_sim_mode?"functional simulation":"performance simulation");
    gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-   kernel_config &config = g_cuda_launch_stack.back();
+   kernel_config config = g_cuda_launch_stack.back();
+   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);
-   if( g_ptx_sim_mode )
-      gpgpu_cuda_ptx_sim_main_func( grid, config.grid_dim(), config.block_dim(), config.get_args() );
-   else
-      gpgpu_cuda_ptx_sim_main_perf( grid, config.grid_dim(), config.block_dim(), config.get_args() );
+   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",
+          kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
+   stream_operation op(grid,g_ptx_sim_mode,stream);
+   g_stream_manager->push(op);
+   g_cuda_launch_stack.pop_back();
    return g_last_cudaError = cudaSuccess;
 }
 
@@ -943,8 +932,10 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
    printf("GPGPU-Sim PTX: cudaStreamCreate\n");
 #if (CUDART_VERSION >= 3000)
    *stream = new struct CUstream_st();
+   g_stream_manager->add_stream(*stream);
 #else
    *stream = 0;
+   printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
 #endif
    return g_last_cudaError = cudaSuccess;
 }
@@ -952,21 +943,33 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 {
 #if (CUDART_VERSION >= 3000)
-   delete stream;
+    g_stream_manager->destroy_stream(stream);
 #endif
    return g_last_cudaError = cudaSuccess;
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
 {
-   printf("GPGPU-Sim PTX: WARNING: This implementation of %s is only a stub! \n", __my_func__);
-   return g_last_cudaError = cudaSuccess; // it is stub because all cuda calls are synchronous
+#if (CUDART_VERSION >= 3000)
+    if( stream == NULL ) 
+        return g_last_cudaError = cudaErrorInvalidResourceHandle;
+    stream->synchronize();
+#else
+    printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
+#endif
+    return g_last_cudaError = cudaSuccess;
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
 {
-   printf("GPGPU-Sim PTX: WARNING: This implementation of %s is only a stub! \n", __my_func__);
+#if (CUDART_VERSION >= 3000)
+   if( stream == NULL ) 
+       return g_last_cudaError = cudaErrorInvalidResourceHandle;
+   return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady;
+#else
+   printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
    return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous
+#endif
 }
 
 /*******************************************************************************
@@ -1005,15 +1008,15 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s
 {
    CUevent_st *e = get_event(event);
    if( !e ) return g_last_cudaError = cudaErrorUnknown;
-   time_t wallclock = time((time_t *)NULL);
-   e->update( gpu_tot_sim_cycle, wallclock );
+   stream_operation op(e,stream);
+   g_stream_manager->push(op);
    return g_last_cudaError = cudaSuccess;
 }
 
 __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)
 {
    CUevent_st *e = get_event(event);
-   if( e == NULL || e->num_updates() == 0 ) {
+   if( e == NULL ) {
       return g_last_cudaError = cudaErrorInvalidValue;
    } else if( e->done() ) {
       return g_last_cudaError = cudaSuccess;
@@ -1024,7 +1027,13 @@ __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)
 
 __host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event)
 {
-   return g_last_cudaError = cudaSuccess;
+    printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n");
+    fflush(stdout);
+    while( !event->done() )
+        ; 
+    printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n");
+    fflush(stdout);
+    return g_last_cudaError = cudaSuccess;
 }
 
 __host__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event)
@@ -1061,22 +1070,20 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start
 
 __host__ cudaError_t CUDARTAPI cudaThreadExit(void)
 {
-   // TODO...  manage memory resources?
+   exit_simulation();
    return g_last_cudaError = cudaSuccess;
 }
 
-
 __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void)
 {
-   //Called on host side 
-   //TODO This function should syncronize if we support Asyn kernel calls
-   return g_last_cudaError = cudaSuccess;
+    //Called on host side 
+    synchronize();
+    return g_last_cudaError = cudaSuccess;
 };
 
 int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
 {
-   //TODO This function should syncronize if we support Asyn kernel calls
-   return g_last_cudaError = cudaSuccess;
+   return cudaThreadExit();
 }
 
 
@@ -1563,12 +1570,6 @@ kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
        printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found\n");
        abort();
    }
-   std::string kname = entry->get_name();
-
-   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;
    for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index c3f5e8a..8943061 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -623,6 +623,7 @@ class _cl_device_id *GPGPUSim_Init()
       gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); 
       the_device = new _cl_device_id(the_gpu);
    } 
+   start_sim_thread(2);
    return the_device;
 }
 
@@ -874,9 +875,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 
    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, GridDim, BlockDim, params );
+      gpgpu_opencl_ptx_sim_main_func( grid );
    else
-      gpgpu_opencl_ptx_sim_main_perf( grid, GridDim, BlockDim, params );
+      gpgpu_opencl_ptx_sim_main_perf( grid );
    return CL_SUCCESS;
 }
 
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 3f2883a..f779e3d 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -1,6 +1,7 @@
 #include "abstract_hardware_model.h"
 #include "cuda-sim/memory.h"
 #include "option_parser.h"
+#include "cuda-sim/ptx_ir.h"
 #include <algorithm>
 
 unsigned mem_access_t::sm_next_access_uid = 0;   
@@ -306,3 +307,11 @@ void warp_inst_t::generate_mem_accesses()
 
     m_mem_accesses_created=true;
 } 
+
+unsigned kernel_info_t::m_next_uid = 1;
+
+
+std::string kernel_info_t::name() const
+{
+    return m_kernel_entry->get_name();
+}
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index c519be1..77a35a7 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -91,6 +91,8 @@ public:
    {
       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 )
    {
@@ -102,8 +104,22 @@ public:
       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++;
    }
 
+   void inc_running() { m_num_cores_running++; }
+   void dec_running()
+   {
+       assert( m_num_cores_running > 0 );
+       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());
+   }
    class function_info *entry() { return m_kernel_entry; }
    const class function_info *entry() const { return m_kernel_entry; }
 
@@ -143,15 +159,22 @@ public:
    {
       return m_next_tid.z < m_block_dim.z && m_next_tid.y < m_block_dim.y && m_next_tid.z < m_block_dim.x;
    }
+   unsigned get_uid() const { return m_uid; }
+   std::string name() const;
 
 private:
    bool m_valid;
    class function_info *m_kernel_entry;
 
+   unsigned m_uid;
+   static unsigned m_next_uid;
+
    dim3 m_grid_dim;
    dim3 m_block_dim;
    dim3 m_next_cta;
    dim3 m_next_tid;
+
+   unsigned m_num_cores_running;
 };
 
 struct core_config {
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 9b65c25..b6a15e8 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1331,6 +1331,8 @@ int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no noti
 
 extern "C" int ptx_debug;
 
+bool g_cuda_launch_blocking = false;
+
 void read_sim_environment_variables() 
 {
    ptx_debug = 0;
@@ -1375,7 +1377,12 @@ void read_sim_environment_variables()
         fflush(stdout);
         g_override_embedded_ptx = true;
     }
+    char *blocking = getenv("CUDA_LAUNCH_BLOCKING");
+    if( blocking && !strcmp(blocking,"1") ) {
+        g_cuda_launch_blocking = true;
+    }
 #else
+   g_cuda_launch_blocking = true;
    g_override_embedded_ptx = true;
 #endif
 
@@ -1388,7 +1395,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, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args)
+void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel )
 {
    printf("GPGPU-Sim: Performing Functional Simulation...\n");
 
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index e7c1965..ef045fb 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -25,7 +25,7 @@ extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *e
                                             struct dim3 gridDim, 
                                             struct dim3 blockDim, 
                                                           class gpgpu_t *gpu );
-extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args);
+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/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 9bb6a23..9ae05d8 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -276,6 +276,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
    option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, 
                   "Clock Domain Frequencies in MhZ {<Core Clock>:<ICNT Clock>:<L2 Clock>:<DRAM Clock>}",
                   "500.0:2000.0:2000.0:2000.0");
+   option_parser_register(opp, "-gpgpu_max_concurrent_kernel", OPT_INT32, &max_concurrent_kernel,
+                          "maximum kernels that can run concurrently on GPU", "8" );
    option_parser_register(opp, "-gpgpu_cflog_interval", OPT_INT32, &gpgpu_cflog_interval, 
                "Interval between each snapshot in control flow logger", 
                "0");
@@ -307,7 +309,6 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound)
    }
 }
 
-
 void gpgpu_sim::launch( kernel_info_t &kinfo )
 {
    unsigned cta_size = kinfo.threads_per_cta();
@@ -319,15 +320,57 @@ void gpgpu_sim::launch( kernel_info_t &kinfo )
       printf("                 modify the CUDA source to decrease the kernel block size.\n");
       abort();
    }
+   unsigned n=0;
+   for(n=0; n < m_running_kernels.size(); n++ ) {
+       if( m_running_kernels[n].done() ) {
+           m_running_kernels[n] = kinfo;
+           break;
+       }
+   }
+   assert(n < m_running_kernels.size());
+}
+
+bool gpgpu_sim::can_start_kernel()
+{
+   for(unsigned n=0; n < m_running_kernels.size(); n++ ) {
+       if( m_running_kernels[n].done() ) 
+           return true;
+   }
+   return false;
+}
 
-   m_running_kernels.push_back(kinfo);
+bool gpgpu_sim::get_more_cta_left() const
+{ 
+   if (m_config.gpu_max_cta_opt != 0) {
+      if( m_total_cta_launched >= m_config.gpu_max_cta_opt )
+          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() ) 
+           return true;
+   }
+   return false;
+}
+
+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() ) {
+            m_last_issued_kernel=idx;
+            return &m_running_kernels[idx];
+        }
+    }
+    return NULL;
 }
 
-kernel_info_t *gpgpu_sim::next_grid()
+unsigned gpgpu_sim::finished_kernel()
 {
-   m_the_kernel = m_running_kernels.front();
-   m_running_kernels.pop_front();
-   return &m_the_kernel;
+    if( m_finished_kernel.empty() ) 
+        return 0;
+    unsigned result = m_finished_kernel.front();
+    m_finished_kernel.pop_front();
+    return result;
 }
 
 void set_ptx_warp_size(const struct core_config * warp_size);
@@ -360,6 +403,10 @@ 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_last_issued_kernel = 0;
+    m_last_cluster_issue = 0;
 }
 
 int gpgpu_sim::shared_mem_size() const
@@ -421,94 +468,76 @@ void gpgpu_sim::reinit_clock_domains(void)
    l2_time = 0;
 }
 
-// return the number of cycle required to run all the trace on the gpu 
-unsigned int gpgpu_sim::run_gpu_sim() 
+bool gpgpu_sim::active()
 {
-   // run a CUDA grid on the GPU microarchitecture simulator
-   kernel_info_t &entry = m_the_kernel;
-   size_t program_size = get_kernel_code_size(entry.entry());
-
-   int not_completed;
-   int mem_busy;
-   int icnt2mem_busy;
-
-   gpu_sim_cycle = 0;
-   not_completed = 1;
-   mem_busy = 1;
-   icnt2mem_busy = 1;
-   more_thread = true;
-   g_total_cta_left=0;
-   gpu_sim_insn = 0;
-   m_shader_stats->new_grid();
+    if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) 
+       return false;
+    if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) 
+       return false;
+    if (m_config.gpu_max_cta_opt && (gpu_tot_issued_cta >= m_config.gpu_max_cta_opt) )
+       return false;
+    if (m_config.gpu_deadlock_detect && gpu_deadlock) 
+       return false;
+    for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) 
+       if( m_cluster[i]->get_not_completed()>0 ) 
+           return true;;
+    for (unsigned i=0;i<m_memory_config->m_n_mem;i++) 
+       if( m_memory_partition_unit[i]->busy()>0 )
+           return true;;
+    if( icnt_busy() )
+        return true;
+    if( get_more_cta_left() )
+        return true;
+    return false;
+}
 
-   reinit_clock_domains();
-   set_param_gpgpu_num_shaders(m_config.num_shader());
-   for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) 
-      m_cluster[i]->reinit();
-   if (m_config.gpu_max_cta_opt != 0) {
-      g_total_cta_left = m_config.gpu_max_cta_opt;
-   } else {
-      g_total_cta_left =  m_the_kernel.num_blocks();
-   }
-   if (m_config.gpu_max_cta_opt != 0) {
-      // the maximum number of CTA has been reached, stop any further simulation
-      if (gpu_tot_issued_cta >= m_config.gpu_max_cta_opt) 
-         return 0;
-   }
+void gpgpu_sim::init()
+{
+    // run a CUDA grid on the GPU microarchitecture simulator
+    gpu_sim_cycle = 0;
+    gpu_sim_insn = 0;
+    m_total_cta_launched=0;
 
-   if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) {
-      return gpu_sim_cycle;
-   }
-   if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) {
-      return gpu_sim_cycle;
-   }
+    reinit_clock_domains();
+    set_param_gpgpu_num_shaders(m_config.num_shader());
+    for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) 
+       m_cluster[i]->reinit();
+    m_shader_stats->new_grid();
+    // initialize the control-flow, memory access, memory latency logger
+    create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval );
+    shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval);
+    if (m_config.gpgpu_cflog_interval != 0) {
+       insn_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size );
+       shader_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, m_config.gpgpu_cflog_interval);
+       shader_mem_acc_create( m_config.num_shader(), m_memory_config->m_n_mem, 4, m_config.gpgpu_cflog_interval);
+       shader_mem_lat_create( m_config.num_shader(), m_config.gpgpu_cflog_interval);
+       shader_cache_access_create( m_config.num_shader(), 3, m_config.gpgpu_cflog_interval);
+       set_spill_interval (m_config.gpgpu_cflog_interval * 40);
+    }
 
-   // initialize the control-flow, memory access, memory latency logger
-   create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, program_size, 0, m_config.gpgpu_cflog_interval );
-   shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval);
-   if (m_config.gpgpu_cflog_interval != 0) {
-      insn_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, program_size );
-      shader_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, m_config.gpgpu_cflog_interval);
-      shader_mem_acc_create( m_config.num_shader(), m_memory_config->m_n_mem, 4, m_config.gpgpu_cflog_interval);
-      shader_mem_lat_create( m_config.num_shader(), m_config.gpgpu_cflog_interval);
-      shader_cache_access_create( m_config.num_shader(), 3, m_config.gpgpu_cflog_interval);
-      set_spill_interval (m_config.gpgpu_cflog_interval * 40);
-   }
+    if (g_network_mode) 
+       icnt_init_grid(); 
+}
 
-   if (g_network_mode) 
-      icnt_init_grid(); 
+void gpgpu_sim::print_stats()
+{
+    m_memory_stats->memlatstat_lat_pw();
+    gpu_tot_sim_cycle += gpu_sim_cycle;
+    gpu_tot_sim_insn += gpu_sim_insn;
 
-   last_gpu_sim_insn = 0;
-   while (not_completed || mem_busy || icnt2mem_busy) {
-      cycle();
-      not_completed = 0;
-      for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) 
-         not_completed += m_cluster[i]->get_not_completed();
-      mem_busy = 0; 
-      for (unsigned i=0;i<m_memory_config->m_n_mem;i++) 
-         mem_busy += m_memory_partition_unit[i]->busy();
-      icnt2mem_busy = icnt_busy();
-      if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) 
-         break;
-      if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) 
-         break;
-      if (m_config.gpu_deadlock_detect && gpu_deadlock) 
-         break;
-   }
-   m_memory_stats->memlatstat_lat_pw();
-   gpu_tot_sim_cycle += gpu_sim_cycle;
-   gpu_tot_sim_insn += gpu_sim_insn;
-   
-   ptx_file_line_stats_write_file();
+    ptx_file_line_stats_write_file();
+    gpu_print_stat();
 
-   gpu_print_stat();
-   if (g_network_mode) {
-      interconnect_stats();
-      printf("----------------------------Interconnect-DETAILS---------------------------------" );
-      icnt_overal_stat();
-      printf("----------------------------END-of-Interconnect-DETAILS-------------------------" );
-   }
+    if (g_network_mode) {
+       interconnect_stats();
+       printf("----------------------------Interconnect-DETAILS---------------------------------" );
+       icnt_overal_stat();
+       printf("----------------------------END-of-Interconnect-DETAILS-------------------------" );
+    }
+}
 
+void gpgpu_sim::deadlock_check()
+{
    if (m_config.gpu_deadlock_detect && gpu_deadlock) {
       fflush(stdout);
       printf("\n\nGPGPU-Sim uArch: ERROR ** deadlock detected: last writeback core %u @ gpu_sim_cycle %u (+ gpu_tot_sim_cycle %u) (%u cycles ago)\n", 
@@ -545,7 +574,6 @@ unsigned int gpgpu_sim::run_gpu_sim()
       fflush(stdout);
       abort();
    }
-   return gpu_sim_cycle;
 }
 
 void gpgpu_sim::gpu_print_stat() const
@@ -634,11 +662,11 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst)
 
 void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) 
 {
+    set_max_cta(kernel);
+
     // find a free CTA context 
     unsigned free_cta_hw_id=(unsigned)-1;
-    unsigned max_concurrent_cta_this_kernel = m_config->max_cta(kernel);
-    assert( max_concurrent_cta_this_kernel <= MAX_CTA_PER_SHADER );
-    for (unsigned i=0;i<max_concurrent_cta_this_kernel;i++ ) {
+    for (unsigned i=0;i<kernel_max_cta_per_shader;i++ ) {
       if( m_cta_status[i]==0 ) {
          free_cta_hw_id=i;
          break;
@@ -722,6 +750,18 @@ int gpgpu_sim::next_clock_domain(void)
    return mask;
 }
 
+void gpgpu_sim::issue_block2core()
+{
+    for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) {
+        unsigned idx = (i+m_last_cluster_issue+1) % m_shader_config->n_simt_clusters;
+        unsigned num = m_cluster[idx]->issue_block2core();
+        if( num ) {
+            m_last_cluster_issue=idx;
+            m_total_cta_launched += num;
+        }
+    }
+}
+
 unsigned long long g_single_step=0; // set this in gdb to single step the pipeline
 
 void gpgpu_sim::cycle()
@@ -778,10 +818,11 @@ void gpgpu_sim::cycle()
       icnt_transfer();
    }
 
+   last_gpu_sim_insn = 0;
    if (clock_mask & CORE) {
       // L1 cache + shader core pipeline stages 
       for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) {
-         if (m_cluster[i]->get_not_completed() || more_thread) {
+         if (m_cluster[i]->get_not_completed() || get_more_cta_left() ) {
                m_cluster[i]->core_cycle();
          }
       }
@@ -791,18 +832,9 @@ void gpgpu_sim::cycle()
       gpu_sim_cycle++;
       if( g_interactive_debugger_enabled ) 
          gpgpu_debug();
-
-      for (unsigned i=0;i<m_shader_config->n_simt_clusters && more_thread;i++) {
-         if ( ( (m_cluster[i]->get_n_active_cta()+1) <= m_cluster[i]->max_cta(m_the_kernel) ) && g_total_cta_left ) {
-            unsigned num = m_cluster[i]->issue_block2core( m_the_kernel );
-            if (num >= g_total_cta_left) {
-               g_total_cta_left = 0;
-               more_thread = false;
-            } else 
-               g_total_cta_left -= num;
-         }
-      }
-
+     
+      issue_block2core();
+      
       // Flush the caches once all of threads are completed.
       if (m_config.gpgpu_flush_cache) {
          int all_threads_complete = 1 ; 
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index db0b18f..5631ab9 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -192,20 +192,8 @@ public:
         m_valid=true;
     }
 
-    void set_max_cta( const kernel_info_t &kernel ) 
-    {
-        // calcaulte the max cta count and cta size for local memory address mapping
-        m_shader_config.gpu_max_cta_per_shader = m_shader_config.max_cta(kernel);
-        //gpu_max_cta_per_shader is limited by number of CTAs if not enough    
-        if( kernel.num_blocks() < m_shader_config.gpu_max_cta_per_shader*num_shader() ) { 
-           m_shader_config.gpu_max_cta_per_shader = (kernel.num_blocks() / num_shader());
-           if (kernel.num_blocks() % num_shader())
-              m_shader_config.gpu_max_cta_per_shader++;
-        }
-        unsigned int gpu_cta_size = kernel.threads_per_cta();
-        m_shader_config.gpu_padded_cta_size = (gpu_cta_size%32) ? 32*((gpu_cta_size/32)+1) : gpu_cta_size;
-    }
     unsigned num_shader() const { return m_shader_config.num_shader(); }
+    unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; }
 
 private:
     void init_clock_domains(void ); 
@@ -234,6 +222,7 @@ private:
     int   gpgpu_dram_sched_queue_size; 
     int   gpgpu_cflog_interval;
     char * gpgpu_clock_domains;
+    unsigned max_concurrent_kernel;
 
     // visualizer
     bool  g_visualizer_enabled;
@@ -254,12 +243,17 @@ public:
    void set_prop( struct cudaDeviceProp *prop );
 
    void launch( kernel_info_t &kinfo );
-   kernel_info_t *next_grid();
+   bool can_start_kernel();
+   unsigned finished_kernel();
+   void set_kernel_done( unsigned uid ) { m_finished_kernel.push_back(uid); }
 
-   unsigned run_gpu_sim();
+   void init();
+   void cycle();
+   bool active(); 
+   void print_stats();
+   void deadlock_check();
 
    void get_pdom_stack_top_info( unsigned sid, unsigned tid, unsigned *pc, unsigned *rpc );
-   const kernel_info_t &the_kernel() const { return m_the_kernel; }
 
    int shared_mem_size() const;
    int num_registers_per_core() const;
@@ -269,6 +263,8 @@ public:
    enum divergence_support_t simd_model() const; 
 
    unsigned threads_per_core() const;
+   bool get_more_cta_left() const;
+   kernel_info_t *select_kernel();
 
    const gpgpu_sim_config &get_config() const { return m_config; }
    void gpu_print_stat() const;
@@ -278,8 +274,8 @@ private:
    // clocks
    void reinit_clock_domains(void);
    int  next_clock_domain(void);
+   void issue_block2core();
     
-   void cycle();
    void L2c_print_cache_stat() const;
    void shader_print_runtime_stat( FILE *fout );
    void shader_print_l1_miss_stat( FILE *fout );
@@ -293,11 +289,12 @@ private:
    class simt_core_cluster **m_cluster;
    class memory_partition_unit **m_memory_partition_unit;
 
-   kernel_info_t m_the_kernel;
-   std::list<kernel_info_t> m_running_kernels;
+   std::vector<kernel_info_t> m_running_kernels;
+   unsigned m_last_issued_kernel;
 
-   unsigned g_total_cta_left;
-   bool more_thread;
+   std::list<unsigned> m_finished_kernel;
+   unsigned m_total_cta_launched;
+   unsigned m_last_cluster_issue;
 
    // time of next rising edge 
    double core_time;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 7a0ad59..07aeb6d 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -110,6 +110,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                   shader_core_stats *stats )
    : m_barriers( config->max_warps_per_shader, config->max_cta_per_core )
 {
+   m_kernel = NULL;
    m_gpu = gpu;
    m_cluster = cluster;
    m_config = config;
@@ -519,7 +520,8 @@ void shader_core_ctx::fetch()
                 for( unsigned t=0; t<m_config->warp_size;t++) {
                     unsigned tid=warp_id*m_config->warp_size+t;
                     if( m_thread[tid].m_functional_model_thread_state ) {
-                        register_cta_thread_exit(tid);
+                        unsigned cta_id = m_warp[warp_id].get_cta_id();
+                        register_cta_thread_exit(cta_id);
                         m_not_completed -= 1;
                         m_active_threads.reset(tid);
                         m_thread[tid].m_functional_model_thread_state=NULL;
@@ -725,26 +727,37 @@ address_type shader_core_ctx::translate_local_memaddr( address_type localaddr, u
 {
    // During functional execution, each thread sees its own memory space for local memory, but these
    // need to be mapped to a shared address space for timing simulation.  We do that mapping here.
-   localaddr /=4;
+
+   address_type thread_base = 0;
+   unsigned max_concurrent_threads=0;
    if (m_config->gpgpu_local_mem_map) {
-      // Dnew = D*nTpC*nCpS*nS + nTpC*C + T%nTpC
-      // C = S + nS*(T/nTpC)
-      // D = data index; T = thread; C = CTA; S = shader core; p = per
-      // keep threads in a warp contiguous
+      // Dnew = D*N + T%nTpC + nTpC*C
+      // N = nTpC*nCpS*nS (max concurent threads)
+      // C = nS*K + S (hw cta number per gpu)
+      // K = T/nTpC   (hw cta number per core)
+      // D = data index
+      // T = thread
+      // nTpC = number of threads per CTA
+      // nCpS = number of CTA per shader
+      // 
+      // for a given local memory address threads in a CTA map to contiguous addresses,
       // then distribute across memory space by CTAs from successive shader cores first, 
       // then by successive CTA in same shader core
-      localaddr *= m_config->gpu_padded_cta_size * m_config->gpu_max_cta_per_shader * num_shader;
-      localaddr += m_config->gpu_padded_cta_size * (m_sid + num_shader * (tid / m_config->gpu_padded_cta_size));
-      localaddr += tid % m_config->gpu_padded_cta_size; 
+      thread_base = 4*(kernel_padded_threads_per_cta * (m_sid + num_shader * (tid / kernel_padded_threads_per_cta))
+                       + tid % kernel_padded_threads_per_cta); 
+      max_concurrent_threads = kernel_padded_threads_per_cta * kernel_max_cta_per_shader * num_shader;
    } else {
       // legacy mapping that maps the same address in the local memory space of all threads 
       // to a single contiguous address region 
-      localaddr *= num_shader * m_config->n_thread_per_shader;
-      localaddr += (m_config->n_thread_per_shader *m_sid) + tid;
+      thread_base = 4*(m_config->n_thread_per_shader * m_sid + tid);
+      max_concurrent_threads = num_shader * m_config->n_thread_per_shader;
    }
-   localaddr *= 4;
+   assert( thread_base < 4/*word size*/*max_concurrent_threads );
 
-   return localaddr;
+   address_type local_word = localaddr/4;
+   address_type word_offset = localaddr%4;
+   address_type linear_address = local_word*max_concurrent_threads + thread_base + word_offset;
+   return linear_address;
 }
 
 /////////////////////////////////////////////////////////////////////////////////////////
@@ -1181,19 +1194,29 @@ void ldst_unit::cycle()
    }
 }
 
-void shader_core_ctx::register_cta_thread_exit(int tid )
+void shader_core_ctx::register_cta_thread_exit( unsigned cta_num )
 {
-   unsigned padded_cta_size = m_gpu->the_kernel().threads_per_cta();
-   if (padded_cta_size%m_config->warp_size) 
-      padded_cta_size = ((padded_cta_size/m_config->warp_size)+1)*(m_config->warp_size);
-   int cta_num = tid/padded_cta_size;
    assert( m_cta_status[cta_num] > 0 );
    m_cta_status[cta_num]--;
    if (!m_cta_status[cta_num]) {
       m_n_active_cta--;
       m_barriers.deallocate_barrier(cta_num);
       shader_CTA_count_unlog(m_sid, 1);
-      printf("GPGPU-Sim uArch: Shader %d finished CTA #%d (%lld,%lld)\n", m_sid, cta_num, gpu_sim_cycle, gpu_tot_sim_cycle );
+      printf("GPGPU-Sim uArch: Shader %d finished CTA #%d (%lld,%lld), %u CTAs running\n", m_sid, cta_num, gpu_sim_cycle, gpu_tot_sim_cycle,
+             m_n_active_cta );
+      if( m_n_active_cta == 0 ) {
+          assert( m_kernel != NULL );
+          m_kernel->dec_running();
+          printf("GPGPU-Sim uArch: Shader %u empty.\n", m_sid );
+          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_kernel=NULL;
+          fflush(stdout);
+      }
    }
 }
 
@@ -1483,11 +1506,20 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const
       printf ("\n");
    }
 
-   if (result < 1) {
-      printf ("GPGPU-Sim uArch: ERROR ** Kernel requires more resources than shader has.\n");
-      abort();
-   }
-   return result;
+    //gpu_max_cta_per_shader is limited by number of CTAs if not enough to keep all cores busy    
+    if( k.num_blocks() < result*num_shader() ) { 
+       result = k.num_blocks() / num_shader();
+       if (k.num_blocks() % num_shader())
+          result++;
+    }
+
+    assert( result <= MAX_CTA_PER_SHADER );
+    if (result < 1) {
+       printf ("GPGPU-Sim uArch: ERROR ** Kernel requires more resources than shader has.\n");
+       abort();
+    }
+
+    return result;
 }
 
 void shader_core_ctx::cycle()
@@ -1712,6 +1744,16 @@ bool shader_core_ctx::warp_waiting_at_mem_barrier( unsigned warp_id )
    return true;
 }
 
+void shader_core_ctx::set_max_cta( const kernel_info_t &kernel ) 
+{
+    // calculate the max cta count and cta size for local memory address mapping
+    kernel_max_cta_per_shader = m_config->max_cta(kernel);
+    unsigned int gpu_cta_size = kernel.threads_per_cta();
+    kernel_padded_threads_per_cta = (gpu_cta_size%m_config->warp_size) ? 
+        m_config->warp_size*((gpu_cta_size/m_config->warp_size)+1) : 
+        gpu_cta_size;
+}
+
 gpgpu_sim *shader_core_ctx::get_gpu()
 {
    return m_gpu;
@@ -2120,13 +2162,21 @@ unsigned simt_core_cluster::get_n_active_cta() const
     return n;
 }
 
-unsigned simt_core_cluster::issue_block2core( class kernel_info_t &kernel )
+unsigned simt_core_cluster::issue_block2core()
 {
     unsigned num_blocks_issued=0;
     for( unsigned i=0; i < m_config->n_simt_cores_per_cluster; i++ ) {
         unsigned core = (i+m_cta_issue_next_core+1)%m_config->n_simt_cores_per_cluster;
-        if( m_core[core]->get_n_active_cta() < m_config->max_cta(kernel) ) {
-            m_core[core]->issue_block2core(kernel);
+        if( m_core[core]->get_not_completed() == 0 ) {
+            if( m_core[core]->get_kernel() == NULL ) {
+                kernel_info_t *k = m_gpu->select_kernel();
+                if( k ) 
+                    m_core[core]->set_kernel(k);
+            }
+        }
+        kernel_info_t *kernel = m_core[core]->get_kernel();
+        if( kernel && !kernel->no_more_ctas_to_run() && (m_core[core]->get_n_active_cta() < m_config->max_cta(*kernel)) ) {
+            m_core[core]->issue_block2core(*kernel);
             num_blocks_issued++;
             m_cta_issue_next_core=core; 
             break;
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 9200c99..7f41cfe 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -983,11 +983,8 @@ struct shader_core_config : public core_config
     unsigned gpgpu_shader_registers;
     int gpgpu_warpdistro_shader;
     unsigned gpgpu_num_reg_banks;
-    unsigned gpu_max_cta_per_shader; // TODO: modify this for fermi... computed based upon kernel 
-                                    // resource usage; used in shader_core_ctx::translate_local_memaddr 
     bool gpgpu_reg_bank_use_warp_id;
     bool gpgpu_local_mem_map;
-    int gpu_padded_cta_size;
     
     unsigned max_sp_latency;
     unsigned max_sfu_latency;
@@ -1122,12 +1119,20 @@ public:
     void cache_flush();
     void accept_fetch_response( mem_fetch *mf );
     void accept_ldst_unit_response( class mem_fetch * mf );
+    void set_kernel( kernel_info_t *k ) 
+    {
+        assert(k);
+        m_kernel=k; 
+        if(k) k->inc_running(); 
+        printf("GPGPU-Sim uArch: Shader %d kernel = 0x%p\n", m_sid, m_kernel );
+    }
     
     // accessors
     bool fetch_unit_response_buffer_full() const;
     bool ldst_unit_response_buffer_full() const;
     unsigned get_not_completed() const { return m_not_completed; }
     unsigned get_n_active_cta() const { return m_n_active_cta; }
+    kernel_info_t *get_kernel() { return m_kernel; }
 
 // used by functional simulation:
     // modifiers
@@ -1148,6 +1153,7 @@ public:
     void dec_inst_in_pipeline( unsigned warp_id ) { m_warp[warp_id].dec_inst_in_pipeline(); } // also used in writeback()
     void store_ack( class mem_fetch *mf );
     bool warp_waiting_at_mem_barrier( unsigned warp_id );
+    void set_max_cta( const kernel_info_t &kernel );
     
     // accessors
     std::list<unsigned> get_regs_written( const inst_t &fvt ) const;
@@ -1163,7 +1169,7 @@ private:
     address_type next_pc( int tid ) const;
     
     void fetch();
-    void register_cta_thread_exit(int cta_num );
+    void register_cta_thread_exit( unsigned cta_num );
     
     void decode();
     void issue_warp( warp_inst_t *&warp, const warp_inst_t *pI, const active_mask_t &active_mask, unsigned warp_id );
@@ -1189,6 +1195,7 @@ private:
     const memory_config *m_memory_config;
     class simt_core_cluster *m_cluster;
     class gpgpu_sim *m_gpu;
+    kernel_info_t *m_kernel;
     
     // statistics 
     shader_core_stats *m_stats;
@@ -1228,6 +1235,10 @@ private:
     ldst_unit *m_ldst_unit;
     static const unsigned MAX_ALU_LATENCY = 64;
     std::bitset<MAX_ALU_LATENCY> m_result_bus;
+
+    // used for local address mapping with single kernel launch
+    unsigned kernel_max_cta_per_shader;
+    unsigned kernel_padded_threads_per_cta;
 };
 
 class simt_core_cluster {
@@ -1243,7 +1254,7 @@ public:
     void icnt_cycle();
 
     void reinit();
-    unsigned issue_block2core( class kernel_info_t &kernel );
+    unsigned issue_block2core();
     void cache_flush();
     bool icnt_injection_buffer_full(unsigned size, bool write);
     void icnt_inject_request_packet(class mem_fetch *mf);
diff --git a/src/gpgpu-sim/stat-tool.cc b/src/gpgpu-sim/stat-tool.cc
index 744a39d..576412c 100644
--- a/src/gpgpu-sim/stat-tool.cc
+++ b/src/gpgpu-sim/stat-tool.cc
@@ -152,7 +152,7 @@ unsigned translate_pc_to_ptxlineno(unsigned pc);
 static int n_thread_CFloggers = 0;
 static thread_CFlocality** thread_CFlogger = NULL;
 
-void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_type start_pc, unsigned long long  logging_interval) 
+void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long  logging_interval) 
 {
    destroy_thread_CFlogger();
    
@@ -163,7 +163,7 @@ void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_t
    char buffer[32];
    for (int i = 0; i < n_thread_CFloggers; i++) {
       snprintf(buffer, 32, "%02d", i);
-      thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, n_insn, start_pc);
+      thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, start_pc);
       if (logging_interval != 0) {
          add_snap_shot_trigger(thread_CFlogger[i]);
          add_spill_log(thread_CFlogger[i]);
@@ -222,10 +222,10 @@ int insn_warp_occ_logger::s_ids = 0;
 
 static std::vector<insn_warp_occ_logger> iwo_logger;
 
-void insn_warp_occ_create( int n_loggers, int simd_width, int n_insn)
+void insn_warp_occ_create( int n_loggers, int simd_width )
 {
    iwo_logger.clear();
-   iwo_logger.assign(n_loggers, insn_warp_occ_logger(simd_width, n_insn));
+   iwo_logger.assign(n_loggers, insn_warp_occ_logger(simd_width));
    for (unsigned i = 0; i < iwo_logger.size(); i++) {
       iwo_logger[i].set_id(i);
    }
@@ -509,12 +509,12 @@ void shader_CTA_count_visualizer_gzprint( gzFile fout )
 ////////////////////////////////////////////////////////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////
 
-thread_insn_span::thread_insn_span(unsigned long long  cycle, int n_insn)
-  : m_cycle(cycle), m_n_insn(n_insn), 
+thread_insn_span::thread_insn_span(unsigned long long  cycle)
+  : m_cycle(cycle),
 #ifdef USE_MAP
      m_insn_span_count() 
 #else 
-     m_insn_span_count(n_insn * 2) 
+     m_insn_span_count(32*1024) 
 #endif
 { 
 }
@@ -522,7 +522,7 @@ thread_insn_span::thread_insn_span(unsigned long long  cycle, int n_insn)
 thread_insn_span::~thread_insn_span() { }
    
 thread_insn_span::thread_insn_span(const thread_insn_span& other)
-      : m_cycle(other.m_cycle), m_n_insn(other.m_n_insn), 
+      : m_cycle(other.m_cycle),
         m_insn_span_count(other.m_insn_span_count) 
 { 
 }
@@ -530,8 +530,7 @@ thread_insn_span::thread_insn_span(const thread_insn_span& other)
 thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other)
 {
    printf("thread_insn_span& operator=\n");
-   if (this != &other && m_n_insn != other.m_n_insn) {
-      m_n_insn = other.m_n_insn;
+   if (this != &other) {
       m_insn_span_count = other.m_insn_span_count;
       m_cycle = other.m_cycle;
    }
@@ -540,7 +539,6 @@ thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other)
    
 thread_insn_span& thread_insn_span::operator+=(const thread_insn_span& other)
 {
-   assert(m_n_insn == other.m_n_insn); // no way to aggregate if they are different programs
    span_count_map::const_iterator i_sc = other.m_insn_span_count.begin();
    for (; i_sc != other.m_insn_span_count.end(); ++i_sc) {
       m_insn_span_count[i_sc->first] += i_sc->second;
@@ -615,12 +613,11 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const
 thread_CFlocality::thread_CFlocality(std::string name, 
                                      unsigned long long  snap_shot_interval, 
                                      int nthreads, 
-                                     int n_insn, 
                                      address_type start_pc, 
                                      unsigned long long  start_cycle)
       : snap_shot_trigger(snap_shot_interval), m_name(name),
         m_nthreads(nthreads), m_thread_pc(nthreads, start_pc), m_cycle(start_cycle),
-        m_thd_span(start_cycle, n_insn)
+        m_thd_span(start_cycle)
 {
    std::fill(m_thread_pc.begin(), m_thread_pc.end(), -1); // so that hw thread with no work assigned will not clobber results
 }
diff --git a/src/gpgpu-sim/stat-tool.h b/src/gpgpu-sim/stat-tool.h
index 00b7771..31a1dcb 100644
--- a/src/gpgpu-sim/stat-tool.h
+++ b/src/gpgpu-sim/stat-tool.h
@@ -118,7 +118,7 @@ public:
 
 class thread_insn_span {
 public:
-   thread_insn_span(unsigned long long  cycle, int n_insn);
+   thread_insn_span(unsigned long long  cycle);
    thread_insn_span(const thread_insn_span& other);
    ~thread_insn_span();
    
@@ -135,14 +135,13 @@ public:
 private: 
    typedef my_hash_map<address_type, int> span_count_map;
    unsigned long long  m_cycle;
-   int m_n_insn;
    span_count_map m_insn_span_count;
 };
 
 class thread_CFlocality : public snap_shot_trigger, public spill_log_interface {
 public:
    thread_CFlocality(std::string name, unsigned long long  snap_shot_interval, 
-                     int nthreads, int n_insn, address_type start_pc, unsigned long long  start_cycle = 0);
+                     int nthreads, address_type start_pc, unsigned long long  start_cycle = 0);
    ~thread_CFlocality();
    
    void update_thread_pc( int thread_id, address_type pc );
@@ -170,9 +169,9 @@ private:
 
 class insn_warp_occ_logger {
 public:
-   insn_warp_occ_logger(int simd_width, int n_insn)
+   insn_warp_occ_logger(int simd_width)
       : m_simd_width(simd_width), 
-        m_insn_warp_occ(n_insn, linear_histogram(1, "", m_simd_width)),
+        m_insn_warp_occ(1,linear_histogram(1, "", m_simd_width)),
         m_id(s_ids++) {}
    
    insn_warp_occ_logger(const insn_warp_occ_logger& other)
@@ -191,7 +190,9 @@ public:
    void set_id(int id) { m_id = id; }
    
    void log(address_type pc, int warp_occ) {
-      m_insn_warp_occ[pc].add2bin(warp_occ - 1);
+       if( pc >= m_insn_warp_occ.size() ) 
+           m_insn_warp_occ.resize(2*pc, linear_histogram(1, "", m_simd_width));
+       m_insn_warp_occ[pc].add2bin(warp_occ - 1);
    }
    
    void print(FILE *fout) const 
@@ -307,7 +308,7 @@ void try_snap_shot (unsigned long long  current_cycle);
 void set_spill_interval (unsigned long long  interval);
 void spill_log_to_file (FILE *fout, int final, unsigned long long  current_cycle);
 
-void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_type start_pc, unsigned long long  logging_interval);
+void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long  logging_interval);
 void destroy_thread_CFlogger( );
 void cflog_update_thread_pc( int logger_id, int thread_id, address_type pc );
 void cflog_snapshot( int logger_id, unsigned long long  cycle );
@@ -316,7 +317,7 @@ void cflog_print_path_expression(FILE *fout);
 void cflog_visualizer_print(FILE *fout);
 void cflog_visualizer_gzprint(gzFile fout);
 
-void insn_warp_occ_create( int n_loggers, int simd_width, int n_insn );
+void insn_warp_occ_create( int n_loggers, int simd_width );
 void insn_warp_occ_log( int logger_id, address_type pc, int warp_occ );
 void insn_warp_occ_print( FILE *fout );
 
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 8579f08..00ba5a2 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -71,6 +71,7 @@
 #include "cuda-sim/ptx_parser.h"
 #include "gpgpu-sim/gpu-sim.h"
 #include "gpgpu-sim/icnt_wrapper.h"
+#include "stream_manager.h"
 
 #include <pthread.h>
 #include <semaphore.h>
@@ -84,30 +85,116 @@ struct gpgpu_ptx_sim_arg *grid_params;
 
 sem_t g_sim_signal_start;
 sem_t g_sim_signal_finish;
+sem_t g_sim_signal_exit;
 time_t g_simulation_starttime;
 pthread_t g_simulation_thread;
 
 gpgpu_sim_config g_the_gpu_config;
 gpgpu_sim *g_the_gpu;
+stream_manager *g_stream_manager;
 
 static void print_simulation_time();
 
-void *gpgpu_sim_thread(void*)
+void *gpgpu_sim_thread_sequential(void*)
 {
-   kernel_info_t *kernel = NULL;
+   // at most one kernel running at a time
+   bool done;
    do {
       sem_wait(&g_sim_signal_start);
-      kernel = g_the_gpu->next_grid();
-      if( kernel ) {
-          g_the_gpu_config.set_max_cta(*kernel);
-          g_the_gpu->run_gpu_sim();
+      done = true;
+      if( g_the_gpu->get_more_cta_left() ) {
+          done = false;
+          g_the_gpu->init();
+          while( g_the_gpu->active() )
+              g_the_gpu->cycle();
+          g_the_gpu->print_stats();
+          g_the_gpu->deadlock_check();
           print_simulation_time();
       }
       sem_post(&g_sim_signal_finish);
-   } while(kernel);
+   } while(!done);
+   sem_post(&g_sim_signal_exit);
    return NULL;
 }
 
+pthread_mutex_t g_sim_lock = PTHREAD_MUTEX_INITIALIZER;
+bool g_sim_active = false;
+bool g_sim_done = true;
+
+void *gpgpu_sim_thread_concurrent(void*)
+{
+    // concurrent kernel execution simulation thread
+    g_the_gpu->init();
+    do {
+        printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n");
+        fflush(stdout);
+        while( g_stream_manager->empty() && !g_sim_done )
+            ;
+        printf("GPGPU-Sim: ** START simulation thread (detected work) **\n");
+        g_stream_manager->print(stdout);
+        fflush(stdout);
+        pthread_mutex_lock(&g_sim_lock);
+        g_sim_active = true;
+        pthread_mutex_unlock(&g_sim_lock);
+        bool active = false;
+        do {
+            // check if a kernel has completed
+            unsigned grid_uid = g_the_gpu->finished_kernel();
+            if( grid_uid )
+                g_stream_manager->register_finished_kernel(grid_uid);
+             
+            // launch operation on device if one is pending and can be run
+            stream_operation op = g_stream_manager->front();
+            op.do_operation(g_the_gpu);
+    
+            // simulate a clock cycle on the GPU 
+            if( g_the_gpu->active() ) 
+                g_the_gpu->cycle();
+            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();
+        pthread_mutex_lock(&g_sim_lock);
+        g_sim_active = false;
+        pthread_mutex_unlock(&g_sim_lock);
+    } while( !g_sim_done );
+    printf("GPGPU-Sim: *** simulation thread exiting ***\n");
+    fflush(stdout);
+    sem_post(&g_sim_signal_exit);
+    return NULL;
+}
+
+void synchronize()
+{
+    printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n");
+    g_stream_manager->print(stdout);
+    fflush(stdout);
+//    sem_wait(&g_sim_signal_finish);
+    bool done = false;
+    do {
+        pthread_mutex_lock(&g_sim_lock);
+        done = g_stream_manager->empty() && !g_sim_active;
+        pthread_mutex_unlock(&g_sim_lock);
+    } while (!done);
+    printf("GPGPU-Sim: detected inactive GPU simulation thread\n");
+    fflush(stdout);
+//    sem_post(&g_sim_signal_start);
+}
+
+void exit_simulation()
+{
+    g_sim_done=true;
+    printf("GPGPU-Sim: exit_simulation called\n");
+    fflush(stdout);
+    sem_wait(&g_sim_signal_exit);
+    printf("GPGPU-Sim: simulation thread signaled exit\n");
+    fflush(stdout);
+}
+
+extern bool g_cuda_launch_blocking;
+
 gpgpu_sim *gpgpu_ptx_sim_init_perf()
 {
    srand(1);
@@ -125,16 +212,29 @@ gpgpu_sim *gpgpu_ptx_sim_init_perf()
    g_the_gpu_config.init();
 
    g_the_gpu = new gpgpu_sim(g_the_gpu_config);
+   g_stream_manager = new stream_manager(g_the_gpu,g_cuda_launch_blocking);
 
    g_simulation_starttime = time((time_t *)NULL);
 
    sem_init(&g_sim_signal_start,0,0);
    sem_init(&g_sim_signal_finish,0,0);
-   pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread,NULL);
+   sem_init(&g_sim_signal_exit,0,0);
 
    return g_the_gpu;
 }
 
+void start_sim_thread(int api)
+{
+    if( g_sim_done ) {
+        g_sim_done = false;
+        if( g_the_gpu_config.get_max_concurrent_kernel() > 1 && api == 1 ) {
+           pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread_concurrent,NULL);
+        } else {
+           pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread_sequential,NULL);
+        }
+    }
+}
+
 void print_simulation_time()
 {
    time_t current_time, difference, d, h, m, s;
@@ -154,21 +254,15 @@ void print_simulation_time()
    fflush(stdout);
 }
 
-int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid,
-                                  struct dim3 gridDim, 
-                                  struct dim3 blockDim, 
-                                  gpgpu_ptx_sim_arg_list_t grid_params )
+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);
+   //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,
-                                  struct dim3 gridDim, 
-                                  struct dim3 blockDim, 
-                                  gpgpu_ptx_sim_arg_list_t grid_params )
+int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid )
 {
    g_the_gpu->launch(grid);
    sem_post(&g_sim_signal_start);
@@ -176,10 +270,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,
-                                  struct dim3 gridDim, 
-                                  struct dim3 blockDim, 
-                                  gpgpu_ptx_sim_arg_list_t grid_params )
+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 689dcf9..b310752 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -71,19 +71,10 @@
 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,
-                                   struct dim3 gridDim, 
-                                   struct dim3 blockDim, 
-                                   gpgpu_ptx_sim_arg_list_t grid_params );
+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,
-                                  struct dim3 gridDim, 
-                                  struct dim3 blockDim, 
-                                  gpgpu_ptx_sim_arg_list_t grid_params );
-
-int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid,
-                                  struct dim3 gridDim, 
-                                  struct dim3 blockDim, 
-                                  gpgpu_ptx_sim_arg_list_t grid_params );
 #endif
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
new file mode 100644
index 0000000..5d1cfe2
--- /dev/null
+++ b/src/stream_manager.cc
@@ -0,0 +1,383 @@
+/*
+ * stream_manager.cc
+ *
+ * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, 
+ * George L. Yuan and the University of British Columbia, Vancouver, 
+ * BC V6T 1Z4, All Rights Reserved.
+ * 
+ * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE
+ * TERMS AND CONDITIONS.
+ * 
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNERS OR CONTRIBUTORS BE
+ * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+ * POSSIBILITY OF SUCH DAMAGE.
+ * 
+ * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h
+ * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda
+ * (property of NVIDIA).  The files benchmarks/BlackScholes/ and 
+ * benchmarks/template/ are derived from the CUDA SDK available from 
+ * http://www.nvidia.com/cuda (also property of NVIDIA).  The files from 
+ * src/intersim/ are derived from Booksim (a simulator provided with the 
+ * textbook "Principles and Practices of Interconnection Networks" available 
+ * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by 
+ * the corresponding legal terms and conditions set forth separately (original 
+ * copyright notices are left in files from these sources and where we have 
+ * modified a file our copyright notice appears before the original copyright 
+ * notice).  
+ * 
+ * Using this version of GPGPU-Sim requires a complete installation of CUDA 
+ * which is distributed seperately by NVIDIA under separate terms and 
+ * conditions.  To use this version of GPGPU-Sim with OpenCL requires a
+ * recent version of NVIDIA's drivers which support OpenCL.
+ * 
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * 
+ * 1. Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ * 
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ * 
+ * 3. Neither the name of the University of British Columbia nor the names of
+ * its contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ * 
+ * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only.  
+ *  
+ * 5. No nonprofit user may place any restrictions on the use of this software,
+ * including as modified by the user, by any other authorized user.
+ * 
+ * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, 
+ * Ali Bakhoda, George L. Yuan, at the University of British Columbia, 
+ * Vancouver, BC V6T 1Z4
+ */
+
+#include "stream_manager.h"
+#include "gpgpusim_entrypoint.h"
+#include "cuda-sim/cuda-sim.h"
+#include "gpgpu-sim/gpu-sim.h"
+
+unsigned CUstream_st::sm_next_stream_uid = 0;
+
+CUstream_st::CUstream_st() 
+{
+    m_pending = false;
+    m_uid = sm_next_stream_uid++;
+    pthread_mutex_init(&m_lock,NULL);
+}
+
+bool CUstream_st::empty()
+{
+    pthread_mutex_lock(&m_lock);
+    bool empty = m_operations.empty();
+    pthread_mutex_unlock(&m_lock);
+    return empty;
+}
+
+bool CUstream_st::busy()
+{
+    pthread_mutex_lock(&m_lock);
+    bool pending = m_pending;
+    pthread_mutex_unlock(&m_lock);
+    return pending;
+}
+
+void CUstream_st::synchronize() 
+{
+    // called by host thread
+    bool done=false;
+    do{
+        pthread_mutex_lock(&m_lock);
+        done = m_operations.empty();
+        pthread_mutex_unlock(&m_lock);
+    } while ( !done );
+}
+
+void CUstream_st::push( const stream_operation &op )
+{
+    // called by host thread
+    pthread_mutex_lock(&m_lock);
+    m_operations.push_back( op );
+    pthread_mutex_unlock(&m_lock);
+}
+
+void CUstream_st::record_next_done()
+{
+    // called by gpu thread
+    pthread_mutex_lock(&m_lock);
+    assert(m_pending);
+    m_operations.pop_front();
+    m_pending=false;
+    pthread_mutex_unlock(&m_lock);
+}
+
+
+stream_operation CUstream_st::next()
+{
+    // called by gpu thread
+    pthread_mutex_lock(&m_lock);
+    m_pending = true;
+    stream_operation result = m_operations.front();
+    pthread_mutex_unlock(&m_lock);
+    return result;
+}
+
+void CUstream_st::print(FILE *fp)
+{
+    pthread_mutex_lock(&m_lock);
+    fprintf(fp,"GPGPU-Sim API:    stream %u has %zu operations\n", m_uid, m_operations.size() );
+    std::list<stream_operation>::iterator i;
+    unsigned n=0;
+    for( i=m_operations.begin(); i!=m_operations.end(); i++ ) {
+        stream_operation &op = *i;
+        fprintf(fp,"GPGPU-Sim API:       %u : ", n++);
+        op.print(fp);
+        fprintf(fp,"\n");
+    }
+    pthread_mutex_unlock(&m_lock);
+}
+
+void stream_operation::do_operation( gpgpu_sim *gpu )
+{
+    if( is_noop() ) 
+        return;
+
+    assert(!m_done && m_stream);
+    printf("GPGPU-Sim API: stream %u performing ", m_stream->get_uid() );
+    switch( m_type ) {
+    case stream_memcpy_host_to_device:
+        printf("memcpy host-to-device\n");
+        gpu->memcpy_to_gpu(m_device_address_dst,m_host_address_src,m_cnt);
+        m_stream->record_next_done();
+        break;
+    case stream_memcpy_device_to_host: 
+        printf("memcpy device-to-host\n");
+        gpu->memcpy_from_gpu(m_host_address_dst,m_device_address_src,m_cnt);
+        m_stream->record_next_done();
+        break;
+    case stream_memcpy_device_to_device:
+        printf("memcpy device-to-device\n");
+        gpu->memcpy_gpu_to_gpu(m_device_address_dst,m_device_address_src,m_cnt); 
+        m_stream->record_next_done();
+        break;
+    case stream_memcpy_to_symbol:
+        printf("memcpy to symbol\n");
+        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
+        m_stream->record_next_done();
+        break;
+    case stream_memcpy_from_symbol:
+        printf("memcpy from symbol\n");
+        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
+        m_stream->record_next_done();
+        break;
+    case stream_kernel_launch:
+        if( gpu->can_start_kernel() ) {
+            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 );
+            else
+                gpu->launch( m_kernel );
+        }
+        break;
+    case stream_event: {
+        printf("event update\n");
+        time_t wallclock = time((time_t *)NULL);
+        m_event->update( gpu_tot_sim_cycle, wallclock );
+        m_stream->record_next_done();
+        } 
+        break;
+    default:
+        abort();
+    }
+    m_done=true;
+    fflush(stdout);
+}
+
+void stream_operation::print( FILE *fp ) const
+{
+    fprintf(fp," stream operation " );
+    switch( m_type ) {
+    case stream_event: fprintf(fp,"event"); break;
+    case stream_kernel_launch: fprintf(fp,"kernel"); break;
+    case stream_memcpy_device_to_device: fprintf(fp,"memcpy device-to-device"); break;
+    case stream_memcpy_device_to_host: fprintf(fp,"memcpy device-to-host"); break;
+    case stream_memcpy_host_to_device: fprintf(fp,"memcpy host-to-device"); break;
+    case stream_memcpy_to_symbol: fprintf(fp,"memcpy to symbol"); break;
+    case stream_memcpy_from_symbol: fprintf(fp,"memcpy from symbol"); break;
+    case stream_no_op: fprintf(fp,"no-op"); break;
+    }
+}
+
+stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) 
+{
+    m_gpu = gpu;
+    m_service_stream_zero = false;
+    m_cuda_launch_blocking = cuda_launch_blocking;
+    pthread_mutex_init(&m_lock,NULL);
+}
+
+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() );
+    stream->record_next_done();
+    m_grid_id_to_stream.erase(grid_uid);
+    pthread_mutex_unlock(&m_lock);
+}
+
+stream_operation stream_manager::front() 
+{
+    // called by gpu simulation thread
+    stream_operation result;
+    pthread_mutex_lock(&m_lock);
+    if( concurrent_streams_empty() )
+        m_service_stream_zero = true;
+    if( m_service_stream_zero ) {
+        if( !m_stream_zero.empty() ) {
+            if( !m_stream_zero.busy() ) {
+                result = m_stream_zero.next();
+                if( result.is_kernel() ) {
+                    unsigned grid_id = result.get_kernel().get_uid();
+                    m_grid_id_to_stream[grid_id] = &m_stream_zero;
+                }
+            }
+        } else {
+            m_service_stream_zero = false;
+        }
+    } else {
+        std::list<struct CUstream_st*>::iterator s;
+        for( s=m_streams.begin(); s != m_streams.end(); s++) {
+            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;
+                    }
+//                }
+                break;
+            }
+        }
+    }
+    pthread_mutex_unlock(&m_lock);
+    return result;
+}
+
+void stream_manager::add_stream( struct CUstream_st *stream )
+{
+    // called by host thread
+    pthread_mutex_lock(&m_lock);
+    m_streams.push_back(stream);
+    pthread_mutex_unlock(&m_lock);
+}
+
+void stream_manager::destroy_stream( CUstream_st *stream )
+{
+    // called by host thread
+    pthread_mutex_lock(&m_lock);
+    while( !stream->empty() )
+        ; 
+    std::list<CUstream_st *>::iterator s;
+    for( s=m_streams.begin(); s != m_streams.end(); s++ ) {
+        if( *s == stream ) {
+            m_streams.erase(s);
+            break;
+        }
+    }
+    delete stream; 
+    pthread_mutex_unlock(&m_lock);
+}
+
+bool stream_manager::concurrent_streams_empty()
+{
+    bool result = true;
+    // called by gpu simulation thread
+    std::list<struct CUstream_st *>::iterator s;
+    for( s=m_streams.begin(); s!=m_streams.end();++s ) {
+        struct CUstream_st *stream = *s;
+        if( !stream->empty() ) {
+            //stream->print(stdout);
+            result = false;
+        }
+    }
+    return result;
+}
+
+bool stream_manager::empty()
+{
+    bool result = true;
+    pthread_mutex_lock(&m_lock);
+    if( !concurrent_streams_empty() ) 
+        result = false;
+    if( !m_stream_zero.empty() ) 
+        result = false;
+    pthread_mutex_unlock(&m_lock);
+    return result;
+}
+
+void stream_manager::print( FILE *fp)
+{
+    pthread_mutex_lock(&m_lock);
+    print_impl(fp);
+    pthread_mutex_unlock(&m_lock);
+}
+
+void stream_manager::print_impl( FILE *fp)
+{
+    fprintf(fp,"GPGPU-Sim API: Stream Manager State\n");
+    std::list<struct CUstream_st *>::iterator s;
+    for( s=m_streams.begin(); s!=m_streams.end();++s ) {
+        struct CUstream_st *stream = *s;
+        if( !stream->empty() ) 
+            stream->print(fp);
+    }
+    if( !m_stream_zero.empty() ) 
+        m_stream_zero.print(fp);
+}
+
+void stream_manager::push( stream_operation op )
+{
+    struct CUstream_st *stream = op.get_stream();
+
+    // block if stream 0 (or concurrency disabled) and pending concurrent operations exist
+    bool block= !stream || m_cuda_launch_blocking;
+    while(block) {
+        pthread_mutex_lock(&m_lock);
+        block = !concurrent_streams_empty();
+        pthread_mutex_unlock(&m_lock);
+    };
+
+    pthread_mutex_lock(&m_lock);
+    if( stream && !m_cuda_launch_blocking ) {
+        stream->push(op);
+    } else {
+        op.set_stream(&m_stream_zero);
+        m_stream_zero.push(op);
+    }
+    print_impl(stdout);
+    pthread_mutex_unlock(&m_lock);
+    if( m_cuda_launch_blocking || stream == NULL ) {
+        while( !empty() )
+            ; 
+    }
+}
+
diff --git a/src/stream_manager.h b/src/stream_manager.h
new file mode 100644
index 0000000..d47f0c7
--- /dev/null
+++ b/src/stream_manager.h
@@ -0,0 +1,294 @@
+/*
+ * stream_manager.h
+ *
+ * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, 
+ * George L. Yuan and the University of British Columbia, Vancouver, 
+ * BC V6T 1Z4, All Rights Reserved.
+ * 
+ * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE
+ * TERMS AND CONDITIONS.
+ * 
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNERS OR CONTRIBUTORS BE
+ * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+ * POSSIBILITY OF SUCH DAMAGE.
+ * 
+ * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h
+ * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda
+ * (property of NVIDIA).  The files benchmarks/BlackScholes/ and 
+ * benchmarks/template/ are derived from the CUDA SDK available from 
+ * http://www.nvidia.com/cuda (also property of NVIDIA).  The files from 
+ * src/intersim/ are derived from Booksim (a simulator provided with the 
+ * textbook "Principles and Practices of Interconnection Networks" available 
+ * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by 
+ * the corresponding legal terms and conditions set forth separately (original 
+ * copyright notices are left in files from these sources and where we have 
+ * modified a file our copyright notice appears before the original copyright 
+ * notice).  
+ * 
+ * Using this version of GPGPU-Sim requires a complete installation of CUDA 
+ * which is distributed seperately by NVIDIA under separate terms and 
+ * conditions.  To use this version of GPGPU-Sim with OpenCL requires a
+ * recent version of NVIDIA's drivers which support OpenCL.
+ * 
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * 
+ * 1. Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ * 
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ * 
+ * 3. Neither the name of the University of British Columbia nor the names of
+ * its contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ * 
+ * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only.  
+ *  
+ * 5. No nonprofit user may place any restrictions on the use of this software,
+ * including as modified by the user, by any other authorized user.
+ * 
+ * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, 
+ * Ali Bakhoda, George L. Yuan, at the University of British Columbia, 
+ * Vancouver, BC V6T 1Z4
+ */
+
+#ifndef STREAM_MANAGER_H_INCLUDED
+#define STREAM_MANAGER_H_INCLUDED
+
+#include "abstract_hardware_model.h"
+#include <list>
+#include <pthread.h>
+#include <time.h>
+
+//class stream_barrier {
+//public:
+//    stream_barrier() { m_pending_streams=0; }
+//    void inc() { m_pending_streams++; }
+//    void dec() { assert(m_pending_streams); m_pending_streams--; }
+//    unsigned value() const { return m_pending_streams; }
+//private:
+//    unsigned m_pending_streams;
+//};
+
+enum stream_operation_type {
+    stream_no_op,
+    stream_memcpy_host_to_device,
+    stream_memcpy_device_to_host,
+    stream_memcpy_device_to_device,
+    stream_memcpy_to_symbol,
+    stream_memcpy_from_symbol,
+    stream_kernel_launch,
+    stream_event
+};
+
+class stream_operation {
+public:
+    stream_operation()
+    {
+        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_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_stream = stream;
+        m_type=stream_memcpy_from_symbol;
+        m_host_address_dst=dst;
+        m_symbol=symbol;
+        m_cnt=count;
+        m_offset=offset;
+        m_done=false;
+    }
+    stream_operation( const kernel_info_t &kernel, bool sim_mode, struct CUstream_st *stream )
+    {
+        m_type=stream_kernel_launch;
+        m_kernel=kernel;
+        m_sim_mode=sim_mode;
+        m_stream=stream;
+        m_done=false;
+    }
+    stream_operation( class CUevent_st *e, struct CUstream_st *stream )
+    {
+        m_type=stream_event;
+        m_event=e;
+        m_stream=stream;
+        m_done=false;
+    }
+    stream_operation( const void *host_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream )
+    {
+        m_type=stream_memcpy_host_to_device;
+        m_host_address_src  =host_address_src;
+        m_device_address_dst=device_address_dst;
+        m_host_address_dst=NULL;
+        m_device_address_src=0;
+        m_cnt=cnt;
+        m_stream=stream;
+        m_sim_mode=false;
+        m_done=false;
+    }
+    stream_operation( size_t device_address_src, void *host_address_dst, size_t cnt, struct CUstream_st *stream  )
+    {
+        m_type=stream_memcpy_device_to_host;
+        m_device_address_src=device_address_src;
+        m_host_address_dst=host_address_dst;
+        m_device_address_dst=0;
+        m_host_address_src=NULL;
+        m_cnt=cnt;
+        m_stream=stream;
+        m_sim_mode=false;
+        m_done=false;
+    }
+    stream_operation( size_t device_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream  )
+    {
+        m_type=stream_memcpy_device_to_device;
+        m_device_address_src=device_address_src;
+        m_device_address_dst=device_address_dst;
+        m_host_address_src=NULL;
+        m_host_address_dst=NULL;
+        m_cnt=cnt;
+        m_stream=stream;
+        m_sim_mode=false;
+        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 ||
+               m_type == stream_memcpy_device_to_host ||
+               m_type == stream_memcpy_host_to_device;
+    }
+    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; }
+    void do_operation( gpgpu_sim *gpu );
+    void print( FILE *fp ) const;
+    struct CUstream_st *get_stream() { return m_stream; }
+    void set_stream( CUstream_st *stream ) { m_stream = stream; }
+
+private:
+    struct CUstream_st *m_stream;
+
+    bool m_done;
+
+    stream_operation_type m_type;
+    size_t      m_device_address_dst;
+    size_t      m_device_address_src;
+    void       *m_host_address_dst;
+    const void *m_host_address_src;
+    size_t      m_cnt;
+
+    const char *m_symbol;
+    size_t m_offset;
+
+    bool m_sim_mode;
+    kernel_info_t m_kernel;
+    class CUevent_st *m_event;
+
+    //stream_barrier *m_barrier;
+};
+
+class CUevent_st {
+public:
+   CUevent_st( bool blocking )
+   {
+      m_uid = ++m_next_event_uid;
+      m_blocking = blocking;
+      m_updates = 0;
+      m_wallclock = 0;
+      m_gpu_tot_sim_cycle = 0;
+      m_done = false;
+   }
+   void update( double cycle, time_t clk )
+   {
+      m_updates++;
+      m_wallclock=clk;
+      m_gpu_tot_sim_cycle=cycle;
+      m_done = true;
+   }
+   //void set_done() { assert(!m_done); m_done=true; }
+   int get_uid() const { return m_uid; }
+   unsigned num_updates() const { return m_updates; }
+   bool done() const { return m_done; }
+   time_t clock() const { return m_wallclock; }
+private:
+   int m_uid;
+   bool m_blocking;
+   bool m_done;
+   int m_updates;
+   time_t m_wallclock;
+   double m_gpu_tot_sim_cycle;
+
+   static int m_next_event_uid;
+};
+
+struct CUstream_st {
+public:
+    CUstream_st(); 
+    bool empty();
+    bool busy();
+    void synchronize();
+    void push( const stream_operation &op );
+    void record_next_done();
+    stream_operation next();
+    stream_operation &front() { return m_operations.front(); }
+    void print( FILE *fp );
+    unsigned get_uid() const { return m_uid; }
+
+private:
+    unsigned m_uid;
+    static unsigned sm_next_stream_uid;
+
+    std::list<stream_operation> m_operations;
+    bool m_pending; // front operation has started but not yet completed
+
+    pthread_mutex_t m_lock; // ensure only one host or gpu manipulates stream operation at one time
+};
+
+class stream_manager {
+public:
+    stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking );
+    void register_finished_kernel( unsigned grid_uid );
+    stream_operation front();
+    void add_stream( CUstream_st *stream );
+    void destroy_stream( CUstream_st *stream );
+    bool concurrent_streams_empty();
+    bool empty();
+    void print( FILE *fp);
+    void push( stream_operation op );
+
+private:
+    void print_impl( FILE *fp);
+
+    bool m_cuda_launch_blocking;
+    gpgpu_sim *m_gpu;
+    std::list<CUstream_st *> m_streams;
+    std::map<unsigned,CUstream_st *> m_grid_id_to_stream;
+    CUstream_st m_stream_zero;
+    bool m_service_stream_zero;
+    pthread_mutex_t m_lock;
+};
+
+#endif