From 6ab2ca48f7f3f97052ae31f56fa93fa7eed84f2c Mon Sep 17 00:00:00 2001 From: WilliamMTK Date: Thu, 6 Mar 2025 14:06:03 -0500 Subject: sst-integration-stream: make SST integration works with streams (#103) * sst-integration-stream: add apis to make SST integration works with stream * Add dev container specs --- src/gpgpusim_entrypoint.cc | 26 ++++++++++++++++++++++++-- 1 file changed, 24 insertions(+), 2 deletions(-) (limited to 'src/gpgpusim_entrypoint.cc') 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& 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); -- cgit v1.3