- 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]

1 files changed, 3 insertions, 2 deletions

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;
 }