Merge pull request #137 from gangmul12/deadlock_fix

Deadlock fix from stream_manager

Author: aamodt

libcuda/cuda_runtime_api.cc

@@ -1746,6 +1746,7 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s
 	if( !e ) return g_last_cudaError = cudaErrorUnknown;
 	struct CUstream_st *s = (struct CUstream_st *)stream;
 	stream_operation op(e,s);
+	e->issue();
 	g_stream_manager->push(op);
 	return g_last_cudaError = cudaSuccess;
 }
@@ -1758,7 +1759,11 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven
    //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html 
 	CUevent_st *e = get_event(event);
 	if( !e ){
-	   printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
+      printf("GPGPU-Sim API: Error at cudaStreamWaitEvent. Event is not created .\n");
+      return g_last_cudaError = cudaErrorInvalidResourceHandle;
+   }
+   else if(e->num_issued() == 0){
+      printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
       return g_last_cudaError = cudaSuccess;
    }
    if (!stream){

src/stream_manager.cc

@@ -194,14 +194,17 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
         m_stream->record_next_done();
         } 
         break;
-    case stream_wait_event: {
+    case stream_wait_event:
         //only allows next op to go if event is done
         //otherwise stays in the stream queue
         printf("stream wait event processing...\n");
-        if(m_event->done())
+        if(m_event->num_updates()>=m_cnt){
             printf("stream wait event done\n");
             m_stream->record_next_done();
         }
+        else{
+            return false;
+        }
         break;
     default:
         abort();
@@ -232,6 +235,7 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking )
     m_service_stream_zero = false;
     m_cuda_launch_blocking = cuda_launch_blocking;
     pthread_mutex_init(&m_lock,NULL);
+    m_last_stream = m_streams.begin();
 }
 
 bool stream_manager::operation( bool * sim)
@@ -330,11 +334,16 @@ stream_operation stream_manager::front()
             m_service_stream_zero = false;
         }
     }
-    
     if(!m_service_stream_zero)
     {
-        std::list<struct CUstream_st*>::iterator s;
-        for( s=m_streams.begin(); s != m_streams.end(); s++) {
+        std::list<struct CUstream_st*>::iterator s = m_last_stream;
+        if(m_last_stream == m_streams.end()){ s = m_streams.begin(); }
+        else{ s++; }
+        for(size_t ii = 0 ; ii < m_streams.size(); ii++, s++) {
+            if(s == m_streams.end()){
+                s = m_streams.begin();
+            }
+            m_last_stream = s;
             CUstream_st *stream = *s;
             if( !stream->busy() && !stream->empty() ) {
                 result = stream->next();
@@ -371,6 +380,7 @@ void stream_manager::destroy_stream( CUstream_st *stream )
         }
     }
     delete stream; 
+    m_last_stream = m_streams.begin();
     pthread_mutex_unlock(&m_lock);
 }
 

src/stream_manager.h

@@ -43,6 +43,45 @@
 //    unsigned m_pending_streams;
 //};
 
+struct 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_issued = 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_updates==m_issued; }
+   time_t clock() const { return m_wallclock; }
+   void issue(){ m_issued++; }
+   unsigned int num_issued() const{ return m_issued; }
+private:
+   int m_uid;
+   bool m_blocking;
+   bool m_done;
+   int m_updates;
+   unsigned int m_issued;
+   time_t m_wallclock;
+   double m_gpu_tot_sim_cycle;
+
+   static int m_next_event_uid;
+};
+
+
 enum stream_operation_type {
     stream_no_op,
     stream_memcpy_host_to_device,
@@ -107,6 +146,7 @@ class stream_operation {
         m_kernel=NULL;
         m_type=stream_wait_event;
         m_event=e;
+        m_cnt = m_event->num_issued();
         m_stream=stream;
         m_done=false;
     }
@@ -163,7 +203,6 @@ class stream_operation {
     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;
 
@@ -183,41 +222,6 @@ class stream_operation {
     kernel_info_t *m_kernel;
     struct CUevent_st *m_event;
 };
-
-struct 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(); 
@@ -268,6 +272,7 @@ class stream_manager {
     CUstream_st m_stream_zero;
     bool m_service_stream_zero;
     pthread_mutex_t m_lock;
+    std::list<struct CUstream_st*>::iterator m_last_stream;
 };
 
 #endif