diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 4 | ||||
| -rw-r--r-- | src/stream_manager.cc | 7 | ||||
| -rw-r--r-- | src/stream_manager.h | 11 |
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; |