MOD: schedule one child kernel each cycle

4 files changed, 146 insertions, 72 deletions

diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index 2a90cba..df2653e 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -20,8 +20,45 @@
       std::cout.flush(); \
    }
 
-std::map<void *, kernel_info_t *> g_cuda_device_launch_map;
-extern stream_manager * g_stream_manager;
+class device_launch_config_t {
+
+public:
+    device_launch_config_t() {}
+
+    device_launch_config_t(dim3 _grid_dim,
+        dim3 _block_dim,
+        unsigned int _shared_mem,
+        function_info * _entry):
+            grid_dim(_grid_dim),
+            block_dim(_block_dim),
+            shared_mem(_shared_mem),
+            entry(_entry) {}
+    
+    dim3 grid_dim;
+    dim3 block_dim;
+    unsigned int shared_mem;
+    function_info * entry;
+
+};
+
+class device_launch_operation_t {
+
+public:
+    device_launch_operation_t() {}
+    device_launch_operation_t(kernel_info_t *_grid,
+        CUstream_st * _stream) :
+            grid(_grid), stream(_stream) {}
+
+    kernel_info_t * grid; //a new child grid
+
+    CUstream_st * stream; 
+
+};
+
+
+std::map<void *, device_launch_config_t> g_cuda_device_launch_param_map;
+std::list<device_launch_operation_t> g_cuda_device_launch_op;
+extern stream_manager *g_stream_manager;
 
 //Handling device runtime api:
 //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize)
@@ -35,8 +72,8 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i
     assert( n_args == 4 );
 
     function_info * child_kernel_entry;
-    struct dim3 gridDim, blockDim;
-    unsigned int sharedMem;
+    struct dim3 grid_dim, block_dim;
+    unsigned int shared_mem;
 
     for( unsigned arg=0; arg < n_args; arg ++ ) {
         const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param#
@@ -48,54 +85,48 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i
 
         if(arg == 0) {//function_info* for the child kernel
             unsigned long long buf;
-			assert(size == sizeof(function_info *));
+            assert(size == sizeof(function_info *));
             thread->m_local_mem->read(from_addr, size, &buf);
             child_kernel_entry = (function_info *)buf;
             assert(child_kernel_entry);
             DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name());
         }
-        else if(arg == 1) { //dim3 gridDim for the child kernel
-			assert(size == sizeof(struct dim3));
-            thread->m_local_mem->read(from_addr, size, & gridDim);
-            DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")");
+        else if(arg == 1) { //dim3 grid_dim for the child kernel
+            assert(size == sizeof(struct dim3));
+            thread->m_local_mem->read(from_addr, size, & grid_dim);
+            DEV_RUNTIME_REPORT("grid (" << grid_dim.x << ", " << grid_dim.y << ", " << grid_dim.z << ")");
         }
-        else if(arg == 2) { //dim3 blockDim for the child kernel
-			assert(size == sizeof(struct dim3));
-            thread->m_local_mem->read(from_addr, size, & blockDim);
-            DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")");
+        else if(arg == 2) { //dim3 block_dim for the child kernel
+            assert(size == sizeof(struct dim3));
+            thread->m_local_mem->read(from_addr, size, & block_dim);
+            DEV_RUNTIME_REPORT("block (" << block_dim.x << ", " << block_dim.y << ", " << block_dim.z << ")");
         }
-        else if(arg == 3) { //unsigned int sharedMem
-			assert(size == sizeof(unsigned int));
-            thread->m_local_mem->read(from_addr, size, & sharedMem);
-            DEV_RUNTIME_REPORT("shared memory " << sharedMem);
+        else if(arg == 3) { //unsigned int shared_mem
+            assert(size == sizeof(unsigned int));
+            thread->m_local_mem->read(from_addr, size, & shared_mem);
+            DEV_RUNTIME_REPORT("shared memory " << shared_mem);
         }
     }
 
-	//get total child kernel argument size and malloc buffer in global memory
-	unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size();
-	void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size);
-	DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer);
-	
-	//create child kernel_info_t and index it with parameter_buffer address
-	kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry);
-    kernel_info_t & parent_grid = thread->get_kernel();
-	DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" <<
-        thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z <<
-        "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z <<
-        ")");
-    child_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid());  
-    assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end());
-    g_cuda_device_launch_map[param_buffer] = child_grid;
+    //get total child kernel argument size and malloc buffer in global memory
+    unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size();
+    void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size);
+    DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer);
 
-	//copy the buffer address to retval0
+    //store param buffer address and launch config
+    device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry);
+    assert(g_cuda_device_launch_param_map.find(param_buffer) == g_cuda_device_launch_param_map.end());
+    g_cuda_device_launch_param_map[param_buffer] = device_launch_config;
+
+    //copy the buffer address to retval0
     const operand_info &actual_return_op = pI->operand_lookup(0); //retval0
     const symbol *formal_return = target_func->get_return_var(); //void *
-	unsigned int return_size = formal_return->get_size_in_bytes();
-	DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size);
-	assert(actual_return_op.is_param_local());
-	assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *));
+    unsigned int return_size = formal_return->get_size_in_bytes();
+    DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size);
+    assert(actual_return_op.is_param_local());
+    assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *));
     addr_t ret_param_addr = actual_return_op.get_symbol()->get_address();
-	thread->m_local_mem->write(ret_param_addr, return_size, &param_buffer, NULL, NULL);
+    thread->m_local_mem->write(ret_param_addr, return_size, &param_buffer, NULL, NULL);
 
 }
 
@@ -109,10 +140,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
     unsigned n_args = target_func->num_args();
     assert( n_args == 2 );
 
-    kernel_info_t * child_grid = NULL;
-    function_info * child_kernel_entry = NULL;
+    kernel_info_t * device_grid = NULL;
+    function_info * device_kernel_entry = NULL;
     void * parameter_buffer;
     struct CUstream_st * child_stream;
+    device_launch_config_t config;
+    device_launch_operation_t device_launch_op;
+
     for( unsigned arg=0; arg < n_args; arg ++ ) {
         const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param#
         const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_#
@@ -123,64 +157,80 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info *
 
         if(arg == 0) {//paramter buffer for child kernel (in global memory)
             //get parameter_buffer from the cudaLaunchDeviceV2_param0
-			assert(size == sizeof(void *));
+            assert(size == sizeof(void *));
             thread->m_local_mem->read(from_addr, size, &parameter_buffer);
             assert((size_t)parameter_buffer >= GLOBAL_HEAP_START);
             DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer);
 
             //get child grid info through parameter_buffer address
-            assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end());
-            child_grid = g_cuda_device_launch_map[parameter_buffer];
-            child_kernel_entry = child_grid->entry();
-            DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name());
+            assert(g_cuda_device_launch_param_map.find(parameter_buffer) != g_cuda_device_launch_param_map.end());
+            config = g_cuda_device_launch_param_map[parameter_buffer];
+            //device_grid = op.grid;
+            device_kernel_entry = config.entry;
+            DEV_RUNTIME_REPORT("find device kernel " << device_kernel_entry->get_name());
+
+            //copy data in parameter_buffer to device kernel param memory
+            unsigned device_kernel_arg_size = device_kernel_entry->get_args_aligned_size();
+            DEV_RUNTIME_REPORT("device_kernel_arg_size " << device_kernel_arg_size);
+            memory_space *device_kernel_param_mem;
 
-            //copy data in parameter_buffer to child kernel param memory
-	        unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size();
-            DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size);
-            memory_space *child_kernel_param_mem = child_grid->get_param_memory();
+            //create child kernel_info_t and index it with parameter_buffer address
+            device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry);
+            kernel_info_t & parent_grid = thread->get_kernel();
+            DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" <<
+                thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z <<
+                "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z <<
+                ")");
+            device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid());  
+            device_launch_op = device_launch_operation_t(device_grid, NULL);
+            device_kernel_param_mem = device_grid->get_param_memory(); //kernel param
+           
             size_t param_start_address = 0;
             //copy in word
-            for(unsigned n = 0; n < child_kernel_arg_size; n += 4) {
+            for(unsigned n = 0; n < device_kernel_arg_size; n += 4) {
                 unsigned int oneword;
                 thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 4, &oneword);
-                child_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); 
+                device_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); 
             }
         }
         else if(arg == 1) { //cudaStream for the child kernel
-			assert(size == sizeof(cudaStream_t));
-            thread->m_local_mem->read(from_addr, size, &child_stream);
 
+            assert(size == sizeof(cudaStream_t));
+            thread->m_local_mem->read(from_addr, size, &child_stream);
+     
             kernel_info_t & parent_kernel = thread->get_kernel();
             if(child_stream == 0) { //default stream on device for current CTA
                 child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); 
-                DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << 
+                DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << 
                     " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream);
             }
             else {
                assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); 
-                DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << 
+                DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << 
                 " to stream " << child_stream->get_uid() << ": " << child_stream);
             }
+    
+                device_launch_op.stream = child_stream;
+            }
         }
-        
     }
 
+  
     //launch child kernel
-    stream_operation op(child_grid, g_ptx_sim_mode, child_stream);
-    g_stream_manager->push(op);
-    g_cuda_device_launch_map.erase(parameter_buffer);
+    g_cuda_device_launch_op.push_back(device_launch_op);
+    g_cuda_device_launch_param_map.erase(parameter_buffer);
 
     //set retval0
     const operand_info &actual_return_op = pI->operand_lookup(0); //retval0
     const symbol *formal_return = target_func->get_return_var(); //cudaError_t
-	unsigned int return_size = formal_return->get_size_in_bytes();
-	DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size);
-	assert(actual_return_op.is_param_local());
-	assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size 
+    unsigned int return_size = formal_return->get_size_in_bytes();
+    DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size);
+    assert(actual_return_op.is_param_local());
+    assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size 
         && return_size == sizeof(cudaError_t));
     cudaError_t error = cudaSuccess;
     addr_t ret_param_addr = actual_return_op.get_symbol()->get_address();
-	thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL);
+    thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL);
 
 }
 
@@ -208,7 +258,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_
         addr_t from_addr = actual_param_op.get_symbol()->get_address();
 
         if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t
-			assert(size == sizeof(cudaStream_t *));
+            assert(size == sizeof(cudaStream_t *));
             thread->m_local_mem->read(from_addr, size, &generic_pStream_addr);
             
             //pStream should be non-zero address in local memory
@@ -217,7 +267,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_
             DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr);
         }
         else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking
-			assert(size == sizeof(unsigned int));
+            assert(size == sizeof(unsigned int));
             thread->m_local_mem->read(from_addr, size, &flags);
             assert(flags == cudaStreamNonBlocking);
         }
@@ -231,13 +281,31 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_
     //set retval0
     const operand_info &actual_return_op = pI->operand_lookup(0); //retval0
     const symbol *formal_return = target_func->get_return_var(); //cudaError_t
-	unsigned int return_size = formal_return->get_size_in_bytes();
-	DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size);
-	assert(actual_return_op.is_param_local());
-	assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size 
+    unsigned int return_size = formal_return->get_size_in_bytes();
+    DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size);
+    assert(actual_return_op.is_param_local());
+    assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size 
         && return_size == sizeof(cudaError_t));
     cudaError_t error = cudaSuccess;
     addr_t ret_param_addr = actual_return_op.get_symbol()->get_address();
-	thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL);
+    thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL);
+
+}
+
 
+void launch_one_device_kernel() {
+    if(!g_cuda_device_launch_op.empty()) {
+        device_launch_operation_t &op = g_cuda_device_launch_op.front();
+
+        stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream);
+        g_stream_manager->push(stream_op);
+        g_cuda_device_launch_op.pop_front();
+    }
+}
+
+void launch_all_device_kernels() {
+    while(!g_cuda_device_launch_op.empty()) {
+        launch_one_device_kernel();
+    }
 }
+
diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h
index 7447d5a..385d605 100644
--- a/src/cuda-sim/cuda_device_runtime.h
+++ b/src/cuda-sim/cuda_device_runtime.h
@@ -6,3 +6,5 @@
 void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
 void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
 void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
+void launch_all_device_kernels();
+void launch_one_device_kernel();
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 47dbc89..126e007 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -61,6 +61,7 @@
 #include "power_stat.h"
 #include "visualizer.h"
 #include "stats.h"
+#include "../cuda-sim/cuda_device_runtime.h"
 
 #ifdef GPGPUSIM_POWER_MODEL
 #include "power_interface.h"
@@ -1370,6 +1371,9 @@ void gpgpu_sim::cycle()
       }
       try_snap_shot(gpu_sim_cycle);
       spill_log_to_file (stdout, 0, gpu_sim_cycle);
+
+      //launch device kernel
+      launch_one_device_kernel();
    }
 }
 
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 546ab3b..f90d9be 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -217,8 +217,8 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking )
 
 bool stream_manager::operation( bool * sim)
 {
-    pthread_mutex_lock(&m_lock);
     bool check=check_finished_kernel();
+    pthread_mutex_lock(&m_lock);
     if(check)m_gpu->print_stats();
     stream_operation op =front();
     if(!op.do_operation( m_gpu )) //not ready to execute