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.cc | 7 +++++++ 1 file changed, 7 insertions(+) (limited to 'src/stream_manager.cc') 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(); } -- cgit v1.3 From 2778712a298a58be4a5abaa28f45805eb3c027c6 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Fri, 25 May 2018 17:41:40 -0700 Subject: seg fault fix --- libcuda/cuda_runtime_api.cc | 16 ++++++++++------ src/stream_manager.cc | 1 + 2 files changed, 11 insertions(+), 6 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index de6d5a5..f130a37 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1157,12 +1157,6 @@ __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) #endif } -__host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags) -{ - stream->push(stream_operation(stream, event, flags)); - return g_last_cudaError = cudaSuccess; -} - /******************************************************************************* * * * * @@ -1205,6 +1199,16 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags) +{ + CUevent_st *e = get_event(event); + if( !e ) return g_last_cudaError = cudaErrorUnknown; + struct CUstream_st *s = (struct CUstream_st *)stream; + stream_operation op(s,e,flags); + g_stream_manager->push(op); + return g_last_cudaError = cudaSuccess; +} + __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event) { CUevent_st *e = get_event(event); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 5b886ab..1b8ebb1 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -199,6 +199,7 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) if(g_debug_execution >= 3) printf("stream wait event\n"); if(m_event->done()) + printf("stream wait event done\n"); m_stream->record_next_done(); } break; -- cgit v1.3 From 89177651e85312fbc841391121908a5e92280881 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Mon, 28 May 2018 14:40:19 -0700 Subject: better print statements for debugging stream_wait_event op --- src/stream_manager.cc | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 1b8ebb1..547e06b 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -188,16 +188,14 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) } break; case stream_event: { - if(g_debug_execution >= 3) - printf("event update\n"); + printf("event update\n"); time_t wallclock = time((time_t *)NULL); m_event->update( gpu_tot_sim_cycle, wallclock ); m_stream->record_next_done(); } break; case stream_wait_event: { - if(g_debug_execution >= 3) - printf("stream wait event\n"); + printf("stream wait event processing...\n"); if(m_event->done()) printf("stream wait event done\n"); m_stream->record_next_done(); -- cgit v1.3 From 21cce426d844c64dac5de7985b97857bb3724e88 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 29 May 2018 10:15:19 -0700 Subject: fixes seg fault on concurrentKernels, speeds up concurrent_streams_empty, adds check for null stream case on cudaStreamWaitEvent --- libcuda/cuda_runtime_api.cc | 4 ++++ src/gpgpusim_entrypoint.cc | 2 +- src/stream_manager.cc | 3 +++ 3 files changed, 8 insertions(+), 1 deletion(-) (limited to 'src/stream_manager.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f130a37..33da17f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1203,6 +1203,10 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven { CUevent_st *e = get_event(event); if( !e ) return g_last_cudaError = cudaErrorUnknown; + if (!stream){ + printf("GPGPU-Sim API: Warning: cudaStreamWaitEvent on NULL stream not currently supported.\n"); + return g_last_cudaError = cudaErrorInvalidValue; + } struct CUstream_st *s = (struct CUstream_st *)stream; stream_operation op(s,e,flags); g_stream_manager->push(op); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index ede9f20..58d8c04 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -108,7 +108,7 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( g_stream_manager->empty() && !g_sim_done ) + while( g_stream_manager->empty_protected() && !g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 547e06b..697d6ec 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -375,6 +375,8 @@ void stream_manager::destroy_stream( CUstream_st *stream ) bool stream_manager::concurrent_streams_empty() { bool result = true; + if (m_streams.empty()) + return true; // called by gpu simulation thread std::list::iterator s; for( s=m_streams.begin(); s!=m_streams.end();++s ) { @@ -382,6 +384,7 @@ bool stream_manager::concurrent_streams_empty() if( !stream->empty() ) { //stream->print(stdout); result = false; + break; } } return result; -- cgit v1.3 From 4b19748d4b3f808d8a45039326673da8caf73186 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 29 May 2018 10:20:39 -0700 Subject: clarifying comments --- src/stream_manager.cc | 2 ++ 1 file changed, 2 insertions(+) (limited to 'src/stream_manager.cc') diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 697d6ec..c6c17ac 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -195,6 +195,8 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) } break; 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()) printf("stream wait event done\n"); -- 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.cc') 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