From fbf8e1180008db2fec31ba1f1e916ea4cb4af97a Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 25 May 2018 17:05:28 -0700 Subject: tentative impl of cudaStreamWaitEvent --- src/stream_manager.h | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) (limited to 'src/stream_manager.h') 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; -- cgit v1.3 From 99a7e032781ee2772cd86e4c1922851f101ec2c7 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 29 May 2018 11:40:51 -0700 Subject: implements null stream case on cudaStreamWaitEvent and updates warning on null event to reflect api reference --- libcuda/cuda_runtime_api.cc | 15 ++++++++------- src/stream_manager.cc | 7 +++++++ src/stream_manager.h | 1 + 3 files changed, 16 insertions(+), 7 deletions(-) (limited to 'src/stream_manager.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index b6fc6c8..dc92522 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1201,18 +1201,19 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags) { + //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: ERROR: Must call cudaEventRecord on event before calling cudaStreamWaitEvent.\n"); - return g_last_cudaError = cudaErrorInvalidValue; + 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){ - printf("GPGPU-Sim API: ERROR: cudaStreamWaitEvent on NULL stream not currently supported.\n"); - return g_last_cudaError = cudaErrorInvalidValue; + g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags); + } else { + struct CUstream_st *s = (struct CUstream_st *)stream; + stream_operation op(s,e,flags); + g_stream_manager->push(op); } - struct CUstream_st *s = (struct CUstream_st *)stream; - stream_operation op(s,e,flags); - g_stream_manager->push(op); return g_last_cudaError = cudaSuccess; } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index c6c17ac..6cd62a2 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -478,3 +478,10 @@ void stream_manager::push( stream_operation op ) } } +void stream_manager::pushCudaStreamWaitEventToAllStreams( CUevent_st *e, unsigned int flags ){ + std::list::iterator s; + for( s=m_streams.begin(); s != m_streams.end(); s++ ) { + stream_operation op(*s,e,flags); + push(op); + } +} diff --git a/src/stream_manager.h b/src/stream_manager.h index edf2b24..1719ebd 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -255,6 +255,7 @@ public: bool empty(); void print( FILE *fp); void push( stream_operation op ); + void pushCudaStreamWaitEventToAllStreams( CUevent_st *e, unsigned int flags ); bool operation(bool * sim); void stop_all_running_kernels(); private: -- cgit v1.3