From 373d64290239f3ed74d98b20494383f03fe189b6 Mon Sep 17 00:00:00 2001 From: Nathan Conrad Date: Sun, 22 Apr 2018 23:31:42 -0400 Subject: Some classes were referred to as a class and a struct (reported as clang warnings). This makes these consistent. --- src/stream_manager.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'src/stream_manager.h') diff --git a/src/stream_manager.h b/src/stream_manager.h index 222a1b2..d3a804f 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -93,7 +93,7 @@ public: m_stream=stream; m_done=false; } - stream_operation( class CUevent_st *e, struct CUstream_st *stream ) + stream_operation( struct CUevent_st *e, struct CUstream_st *stream ) { m_kernel=NULL; m_type=stream_event; @@ -172,10 +172,10 @@ private: bool m_sim_mode; kernel_info_t *m_kernel; - class CUevent_st *m_event; + struct CUevent_st *m_event; }; -class CUevent_st { +struct CUevent_st { public: CUevent_st( bool blocking ) { -- cgit v1.3 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 --- libcuda/cuda_runtime_api.cc | 4 ++-- src/stream_manager.cc | 7 +++++++ src/stream_manager.h | 11 ++++++++++- 3 files changed, 19 insertions(+), 3 deletions(-) (limited to 'src/stream_manager.h') 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; -- 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 From fe58efe9c8ca38f7d0f3781e54b04bc526bdfd07 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 10 Jun 2019 18:06:42 -0400 Subject: fixing thrust error --- libcuda/cuda_runtime_api.cc | 3 +++ src/stream_manager.h | 2 ++ 2 files changed, 5 insertions(+) (limited to 'src/stream_manager.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index fb3e07a..17a5c96 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1498,6 +1498,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = g_cuda_launch_stack.back(); struct CUstream_st *stream = config.get_stream(); + if(g_stream_manager->is_blocking()) + stream = 0; + printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); diff --git a/src/stream_manager.h b/src/stream_manager.h index 91d1b36..3fbdbaf 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -258,6 +258,8 @@ public: void pushCudaStreamWaitEventToAllStreams( CUevent_st *e, unsigned int flags ); bool operation(bool * sim); void stop_all_running_kernels(); + unsigned size() {return m_streams.size(); }; + bool is_blocking() {return m_cuda_launch_blocking; }; private: void print_impl( FILE *fp); -- cgit v1.3