diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/gpgpusim_entrypoint.cc | 26 | ||||
| -rw-r--r-- | src/stream_manager.cc | 40 | ||||
| -rw-r--r-- | src/stream_manager.h | 22 |
3 files changed, 79 insertions, 9 deletions
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index d4c0452..be49229 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -56,7 +56,9 @@ class stream_manager *g_stream_manager() { // SST callback extern void SST_callback_cudaThreadSynchronize_done(); +extern void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream); __attribute__((weak)) void SST_callback_cudaThreadSynchronize_done() {} +__attribute__((weak)) void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream) {} void *gpgpu_sim_thread_sequential(void *ctx_ptr) { gpgpu_context *ctx = (gpgpu_context *)ctx_ptr; @@ -189,12 +191,33 @@ bool SST_Cycle() { // Check if Synchronize is done when SST previously requested // cudaThreadSynchronize if (GPGPU_Context()->requested_synchronize && - ((g_stream_manager()->empty() && !GPGPUsim_ctx_ptr()->g_sim_active) || + ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) || GPGPUsim_ctx_ptr()->g_sim_done)) { SST_callback_cudaThreadSynchronize_done(); GPGPU_Context()->requested_synchronize = false; } + // Polling to check for each stream if it is marked for requested with sync + if (g_stream_manager()->get_stream_zero()->requested_synchronize() && + ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) || + GPGPUsim_ctx_ptr()->g_sim_done)) { + SST_callback_cudaStreamSynchronize_done(0); + g_stream_manager()->get_stream_zero()->reset_request_synchronize(); + } + + // Iterate through each stream to check if SST is waiting on + // it and it does not have any operation + std::list<CUstream_st *>& streams = g_stream_manager()->get_concurrent_streams(); + for (auto it = streams.begin(); it != streams.end(); it++) { + CUstream_st *stream = *it; + if (stream->requested_synchronize() && + stream->empty()) { + // This stream is ready + SST_callback_cudaStreamSynchronize_done(stream); + stream->reset_request_synchronize(); + } + } + if (g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_done && !g_the_gpu()->active()) { GPGPUsim_ctx_ptr()->g_sim_active = false; @@ -272,7 +295,6 @@ void gpgpu_context::synchronize() { bool gpgpu_context::synchronize_check() { // printf("GPGPU-Sim: synchronize checking for inactive GPU simulation\n"); - requested_synchronize = true; the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); // sem_wait(&g_sim_signal_finish); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index bb95bb1..d43964a 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -34,15 +34,20 @@ unsigned CUstream_st::sm_next_stream_uid = 0; -// SST memcpy callbacks -extern void SST_callback_memcpy_H2D_done(); -extern void SST_callback_memcpy_D2H_done(); +// SST memcpy callbacks, called after a stream operation is done via record_next_done() +extern void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream); +extern void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream); extern void SST_callback_memcpy_to_symbol_done(); extern void SST_callback_memcpy_from_symbol_done(); -__attribute__((weak)) void SST_callback_memcpy_H2D_done() {} -__attribute__((weak)) void SST_callback_memcpy_D2H_done() {} +extern void SST_callback_cudaEventSynchronize_done(cudaEvent_t event); +extern void SST_callback_kernel_done(cudaStream_t stream); +__attribute__((weak)) void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {} +__attribute__((weak)) void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {} __attribute__((weak)) void SST_callback_memcpy_to_symbol_done() {} __attribute__((weak)) void SST_callback_memcpy_from_symbol_done() {} +__attribute__((weak)) void SST_callback_cudaEventSynchronize_done(cudaEvent_t event); +__attribute__((weak)) void SST_callback_kernel_done(cudaStream_t stream); + CUstream_st::CUstream_st() { m_pending = false; @@ -74,6 +79,10 @@ void CUstream_st::synchronize() { } while (!done); } +bool CUstream_st::synchronize_check() { + return m_operations.empty(); +} + void CUstream_st::push(const stream_operation &op) { // called by host thread pthread_mutex_lock(&m_lock); @@ -132,13 +141,15 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) { if (g_debug_execution >= 3) printf("memcpy host-to-device\n"); gpu->memcpy_to_gpu(m_device_address_dst, m_host_address_src, m_cnt); m_stream->record_next_done(); - if (gpu->is_SST_mode()) SST_callback_memcpy_H2D_done(); + if (gpu->is_SST_mode()) { + SST_callback_memcpy_H2D_done((uint64_t) m_device_address_dst, (uint64_t) m_host_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream); + } break; case stream_memcpy_device_to_host: if (g_debug_execution >= 3) printf("memcpy device-to-host\n"); gpu->memcpy_from_gpu(m_host_address_dst, m_device_address_src, m_cnt); m_stream->record_next_done(); - if (gpu->is_SST_mode()) SST_callback_memcpy_D2H_done(); + if (gpu->is_SST_mode()) SST_callback_memcpy_D2H_done((uint64_t) m_host_address_dst, (uint64_t) m_device_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream); break; case stream_memcpy_device_to_device: if (g_debug_execution >= 3) printf("memcpy device-to-device\n"); @@ -194,6 +205,13 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) { time_t wallclock = time((time_t *)NULL); m_event->update(gpu->gpu_tot_sim_cycle, wallclock); m_stream->record_next_done(); + if ((gpu->is_SST_mode()) && m_event->done() && + m_event->requested_synchronize()) { + // Notify that the event is done + SST_callback_cudaEventSynchronize_done(m_event); + // Reset the sync flag as we have notified SST + m_event->reset_request_synchronize(); + } } break; case stream_wait_event: // only allows next op to go if event is done @@ -252,6 +270,9 @@ stream_manager::stream_manager(gpgpu_sim *gpu, bool cuda_launch_blocking) { m_cuda_launch_blocking = cuda_launch_blocking; pthread_mutex_init(&m_lock, NULL); m_last_stream = m_streams.begin(); + + // Mark stream zero as the default stream + m_stream_zero.set_stream_zero(); } bool stream_manager::operation(bool *sim) { @@ -303,6 +324,11 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) { // grid_uid, stream->get_uid()); kernel_stat.flush(); // kernel_stat.close(); stream->record_next_done(); + // Callback to notify a kernel is done for SST's stream + // manager to support with nonblocking + blocking kernel launch + if (m_gpu->is_SST_mode()) { + SST_callback_kernel_done(stream->is_stream_zero_stream() ? 0 : stream); + } m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); delete kernel; diff --git a/src/stream_manager.h b/src/stream_manager.h index c1b76d0..55cfb8d 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -69,6 +69,11 @@ struct CUevent_st { void issue() { m_issued++; } unsigned int num_issued() const { return m_issued; } + // SST related, stating this event is requested to synchronize + void set_request_synchronize() { m_requested_synchronize = true; } + void reset_request_synchronize() { m_requested_synchronize = false; } + bool requested_synchronize() const { return m_requested_synchronize; } + private: int m_uid; bool m_blocking; @@ -77,6 +82,9 @@ struct CUevent_st { unsigned int m_issued; time_t m_wallclock; double m_gpu_tot_sim_cycle; + + // SST related + bool m_requested_synchronize = false; static int m_next_event_uid; }; @@ -226,6 +234,7 @@ struct CUstream_st { bool empty(); bool busy(); void synchronize(); + bool synchronize_check(); void push(const stream_operation &op); void record_next_done(); stream_operation next(); @@ -233,6 +242,12 @@ struct CUstream_st { stream_operation &front() { return m_operations.front(); } void print(FILE *fp); unsigned get_uid() const { return m_uid; } + void set_request_synchronize() { m_requested_synchronize = true; } + void reset_request_synchronize() { m_requested_synchronize = false; } + bool requested_synchronize() const { return m_requested_synchronize; } + void set_stream_zero() { is_stream_zero = true; } + bool is_stream_zero_stream() { return is_stream_zero; } + void reset_stream_zero() { is_stream_zero = false; } private: unsigned m_uid; @@ -243,6 +258,11 @@ struct CUstream_st { pthread_mutex_t m_lock; // ensure only one host or gpu manipulates stream // operation at one time + + // SST related, use to record the stream is requested to synchronize + bool m_requested_synchronize = false; + // Whether this is the default stream + bool is_stream_zero = false; }; class stream_manager { @@ -263,6 +283,8 @@ class stream_manager { void stop_all_running_kernels(); unsigned size() { return m_streams.size(); }; bool is_blocking() { return m_cuda_launch_blocking; }; + CUstream_st *get_stream_zero() { return &m_stream_zero; }; + std::list<CUstream_st *>& get_concurrent_streams() { return m_streams; }; private: void print_impl(FILE *fp); |