Remove g_stream_manager()

Signed-off-by: Mengchi Zhang <[email protected]>

1 files changed, 218 insertions, 130 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 8930d61..9459b64 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -176,8 +176,6 @@ struct cudaArray
 
 cudaError_t g_last_cudaError = cudaSuccess;
 
-//extern stream_manager *g_stream_manager();
-
 void register_ptx_function( const char *name, function_info *impl )
 {
 	// no longer need this
@@ -941,7 +939,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
 			kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
 	stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream);
-	g_stream_manager()->push(op);
+	ctx->the_gpgpusim->g_stream_manager->push(op);
 	ctx->api->g_cuda_launch_stack.pop_back();
 	return g_last_cudaError = cudaSuccess;
 }
@@ -1080,6 +1078,50 @@ __host__ cudaError_t CUDARTAPI cudaMallocArrayInternal(struct cudaArray **array,
 	}
 }
 
+__host__ cudaError_t CUDARTAPI cudaMemcpyInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//CUctx_st *context = GPGPUSim_Context();
+	//gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	if(g_debug_execution >= 3)
+		printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
+	if( kind == cudaMemcpyHostToDevice )
+		ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+	else if( kind == cudaMemcpyDeviceToHost )
+		ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
+	else if( kind == cudaMemcpyDeviceToDevice )
+		ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
+	else if ( kind == cudaMemcpyDefault ) {
+		if ((size_t)src >= GLOBAL_HEAP_START) {
+			if ((size_t)dst >= GLOBAL_HEAP_START)
+				ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
+			else
+				ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host
+		}
+		else {
+			if ((size_t)dst >= GLOBAL_HEAP_START)
+				ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+			else {
+				printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n");
+				abort();
+			}
+		}
+	}
+	else {
+		printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
+		abort();
+	}
+	return g_last_cudaError = cudaSuccess;
+}
+
 __host__ cudaError_t CUDARTAPI cudaMemcpyToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
 { 
     gpgpu_context *ctx;
@@ -1174,6 +1216,69 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayInternal(struct cudaArray *dst
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbolInternal(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice), gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//CUctx_st *context = GPGPUSim_Context();
+	assert(kind == cudaMemcpyHostToDevice);
+	printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
+	//stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
+	ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
+	//gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
+	return g_last_cudaError = cudaSuccess;
+}
+
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbolInternal(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost), gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//CUctx_st *context = GPGPUSim_Context();
+	assert(kind == cudaMemcpyDeviceToHost);
+	printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
+	ctx->the_gpgpusim->g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) );
+	//gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyAsyncInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	struct CUstream_st *s = (struct CUstream_st *)stream;
+	switch( kind ) {
+	case cudaMemcpyHostToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
+	case cudaMemcpyDeviceToHost: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
+	case cudaMemcpyDeviceToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
+	default:
+		abort();
+	}
+	return g_last_cudaError = cudaSuccess;
+}
+
+
 #if (CUDART_VERSION >= 8000)
 cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
 {
@@ -1729,6 +1834,48 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernelInternal( const char* hostFun, di
     return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaStreamCreateInternal(cudaStream_t *stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("GPGPU-Sim PTX: cudaStreamCreate\n");
+#if (CUDART_VERSION >= 3000)
+	*stream = new struct CUstream_st();
+	ctx->the_gpgpusim->g_stream_manager->add_stream(*stream);
+#else
+	*stream = 0;
+	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
+#endif
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#if (CUDART_VERSION >= 3000)
+	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
+	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
+	stream->synchronize();
+	ctx->the_gpgpusim->g_stream_manager->destroy_stream(stream);
+#endif
+	return g_last_cudaError = cudaSuccess;
+}
+
 void __cudaRegisterTextureInternal(
 		void **fatCubinHandle,
 		const struct textureReference *hostVar,
@@ -1858,6 +2005,66 @@ CUresult CUDAAPI cuLaunchKernelInternal(CUfunction f,
 }
 #endif /* CUDART_VERSION >= 4000 */
 
+CUevent_st *get_event(cudaEvent_t event)
+{
+	unsigned event_uid;
+#if CUDART_VERSION >= 3000
+	event_uid = event->get_uid();
+#else
+	event_uid = event;
+#endif
+	event_tracker_t::iterator e = g_timer_events.find(event_uid);
+	if( e == g_timer_events.end() )
+		return NULL;
+	return e->second;
+}
+
+__host__ cudaError_t CUDARTAPI cudaEventRecordInternal(cudaEvent_t event, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUevent_st *e = get_event(event);
+	if( !e ) return g_last_cudaError = cudaErrorUnknown;
+	struct CUstream_st *s = (struct CUstream_st *)stream;
+	stream_operation op(e,s);
+	ctx->the_gpgpusim->g_stream_manager->push(op);
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, cudaEvent_t event, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+   //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: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
+      return g_last_cudaError = cudaSuccess;
+   }
+   if (!stream){
+     ctx->the_gpgpusim->g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags);
+   } else {
+	   struct CUstream_st *s = (struct CUstream_st *)stream;
+	   stream_operation op(s,e,flags);
+	   ctx->the_gpgpusim->g_stream_manager->push(op);
+   }
+	return g_last_cudaError = cudaSuccess;
+}
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1930,40 +2137,7 @@ __host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
 
 __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//CUctx_st *context = GPGPUSim_Context();
-	//gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	if(g_debug_execution >= 3)
-		printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
-	if( kind == cudaMemcpyHostToDevice )
-		g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) );
-	else if( kind == cudaMemcpyDeviceToHost )
-		g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) );
-	else if( kind == cudaMemcpyDeviceToDevice )
-		g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) );
-	else if ( kind == cudaMemcpyDefault ) {
-		if ((size_t)src >= GLOBAL_HEAP_START) {
-			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
-			else
-				g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); // device to host
-		}
-		else {
-			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) );
-			else {
-				printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n");
-				abort();
-			}
-		}
-	}
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemcpyInternal(dst, src, count, kind);
 }
 
 __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
@@ -2022,30 +2196,13 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, s
 
 __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//CUctx_st *context = GPGPUSim_Context();
-	assert(kind == cudaMemcpyHostToDevice);
-	printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
-	//stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
-	g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) );
-	//gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemcpyToSymbolInternal(symbol, src, count, offset, kind);
 }
 
 
 __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost))
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//CUctx_st *context = GPGPUSim_Context();
-	assert(kind == cudaMemcpyDeviceToHost);
-	printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
-	g_stream_manager()->push( stream_operation(symbol,dst,count,offset,0) );
-	//gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemcpyFromSymbolInternal(dst, symbol, count, offset, kind);
 }
 
 __host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){
@@ -2067,18 +2224,7 @@ __host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){
 
 __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	switch( kind ) {
-	case cudaMemcpyHostToDevice: g_stream_manager()->push( stream_operation(src,(size_t)dst,count,s) ); break;
-	case cudaMemcpyDeviceToHost: g_stream_manager()->push( stream_operation((size_t)src,dst,count,s) ); break;
-	case cudaMemcpyDeviceToDevice: g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
-	default:
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemcpyAsyncInternal(dst, src, count, kind, stream);
 }
 
 
@@ -2414,18 +2560,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel( const char* hostFun, dim3 gridD
 
 __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	printf("GPGPU-Sim PTX: cudaStreamCreate\n");
-#if (CUDART_VERSION >= 3000)
-	*stream = new struct CUstream_st();
-	g_stream_manager()->add_stream(*stream);
-#else
-	*stream = 0;
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
-#endif
-	return g_last_cudaError = cudaSuccess;
+    return cudaStreamCreateInternal(stream);
 }
 
 //TODO: introduce priorities
@@ -2452,16 +2587,7 @@ __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t
 
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#if (CUDART_VERSION >= 3000)
-	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
-	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
-	stream->synchronize();
-	g_stream_manager()->destroy_stream(stream);
-#endif
-	return g_last_cudaError = cudaSuccess;
+    return cudaStreamDestroyInternal(stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
@@ -2516,52 +2642,14 @@ __host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event)
 	return g_last_cudaError = cudaSuccess;
 }
 
-CUevent_st *get_event(cudaEvent_t event)
-{
-	unsigned event_uid;
-#if CUDART_VERSION >= 3000
-	event_uid = event->get_uid();
-#else
-	event_uid = event;
-#endif
-	event_tracker_t::iterator e = g_timer_events.find(event_uid);
-	if( e == g_timer_events.end() )
-		return NULL;
-	return e->second;
-}
-
 __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUevent_st *e = get_event(event);
-	if( !e ) return g_last_cudaError = cudaErrorUnknown;
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	stream_operation op(e,s);
-	g_stream_manager()->push(op);
-	return g_last_cudaError = cudaSuccess;
+    return cudaEventRecordInternal(event, stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-   //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: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
-      return g_last_cudaError = cudaSuccess;
-   }
-   if (!stream){
-      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);
-   }
-	return g_last_cudaError = cudaSuccess;
+    return cudaStreamWaitEventInternal(stream, event, flags);
 }
 
 __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)