tentative impl of cudaStreamWaitEvent

3 files changed, 19 insertions, 3 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index cfa9cec..de6d5a5 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1159,8 +1159,8 @@ __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
 
 __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags)
 {
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
+   stream->push(stream_operation(stream, event, flags));
+	return g_last_cudaError = cudaSuccess;
 }
 
 /*******************************************************************************
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index 3b6cbd5..5b886ab 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -195,6 +195,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
         m_stream->record_next_done();
         } 
         break;
+    case stream_wait_event: {
+        if(g_debug_execution >= 3)
+            printf("stream wait event\n");
+        if(m_event->done())
+            m_stream->record_next_done();
+        }
+        break;
     default:
         abort();
     }
diff --git a/src/stream_manager.h b/src/stream_manager.h
index 222a1b2..edf2b24 100644
--- a/src/stream_manager.h
+++ b/src/stream_manager.h
@@ -51,7 +51,8 @@ enum stream_operation_type {
     stream_memcpy_to_symbol,
     stream_memcpy_from_symbol,
     stream_kernel_launch,
-    stream_event
+    stream_event,
+    stream_wait_event
 };
 
 class stream_operation {
@@ -101,6 +102,14 @@ public:
         m_stream=stream;
         m_done=false;
     }
+    stream_operation( struct CUstream_st *stream, class CUevent_st *e, unsigned int flags )
+    {
+        m_kernel=NULL;
+        m_type=stream_wait_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_kernel=NULL;