diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 15 | ||||
| -rw-r--r-- | src/stream_manager.cc | 7 | ||||
| -rw-r--r-- | src/stream_manager.h | 1 |
3 files changed, 16 insertions, 7 deletions
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<CUstream_st *>::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: |